diff --git a/clang/lib/AST/Expr.cpp b/clang/lib/AST/Expr.cpp
index f5d1c5a2261dc..41e36393e026a 100644
--- a/clang/lib/AST/Expr.cpp
+++ b/clang/lib/AST/Expr.cpp
@@ -270,7 +270,14 @@ bool Expr::isFlexibleArrayMemberLike(
           continue;
         }
         if (ConstantArrayTypeLoc CTL = TL.getAs<ConstantArrayTypeLoc>()) {
-          const Expr *SizeExpr = dyn_cast<IntegerLiteral>(CTL.getSizeExpr());
+          // FIXME: changed dyn_cast to dyn_cast_or_null
+          //    to work around the fact that CTL.getSizeExpr() isn't set
+          //    for a FieldDecl of a class generated from a lambda capture.
+          //    This is highlighted only by the way lambda expression used
+          //    as a SYCL kernel is being processed.
+          //    In normal situation the capture list is used.
+          //    No harm done, just a work around.
+          const Expr *SizeExpr = dyn_cast_or_null<IntegerLiteral>(CTL.getSizeExpr());
           if (!SizeExpr || SizeExpr->getExprLoc().isMacroID())
             return false;
         }
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 326d3750542b3..fbb507ca40d6c 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -889,10 +889,14 @@ class KernelBodyTransform : public TreeTransform<KernelBodyTransform> {
     auto Ref = dyn_cast<DeclaratorDecl>(DRE->getDecl());
     if (Ref && Ref == MappingPair.first) {
       auto NewDecl = MappingPair.second;
+      QualType ExprType = NewDecl->getType();
+      if (ExprType->isReferenceType()) {
+        ExprType = ExprType->getPointeeType();
+      }
       return DeclRefExpr::Create(
           SemaRef.getASTContext(), DRE->getQualifierLoc(),
           DRE->getTemplateKeywordLoc(), NewDecl, false, DRE->getNameInfo(),
-          NewDecl->getType(), DRE->getValueKind());
+          ExprType, DRE->getValueKind());
     }
     return DRE;
   }
@@ -2800,7 +2804,6 @@ static bool isESIMDKernelType(const CXXRecordDecl *KernelObjType) {
 class SyclKernelBodyCreator : public SyclKernelFieldHandler {
   SyclKernelDeclCreator &DeclCreator;
   llvm::SmallVector<Stmt *, 16> BodyStmts;
-  llvm::SmallVector<InitListExpr *, 16> CollectionInitExprs;
   llvm::SmallVector<Stmt *, 16> FinalizeStmts;
   // This collection contains the information required to add/remove information
   // about arrays as we enter them.  The InitializedEntity component is
@@ -2808,6 +2811,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
   // current element being worked on, which is updated every time we visit
   // nextElement.
   llvm::SmallVector<std::pair<InitializedEntity, uint64_t>, 8> ArrayInfos;
+  const CXXRecordDecl *WrappingUnion;
   VarDecl *KernelObjClone;
   InitializedEntity VarEntity;
   const CXXRecordDecl *KernelObj;
@@ -2839,15 +2843,22 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
     // Push the Kernel function scope to ensure the scope isn't empty
     SemaRef.PushFunctionScope();
 
-    // Initialize kernel object local clone
-    assert(CollectionInitExprs.size() == 1 &&
-           "Should have been popped down to just the first one");
-    KernelObjClone->setInit(CollectionInitExprs.back());
+    // Create a reference to the Functor in the union
+    FieldDecl *WrappedField = *WrappingUnion->field_begin();
+    VarDecl *KernelObjRef = VarDecl::Create(
+        SemaRef.Context, DeclCreator.getKernelDecl(), SourceLocation(),
+        SourceLocation(), WrappedField->getIdentifier(),
+        SemaRef.Context.getLValueReferenceType(WrappedField->getType()),
+        nullptr, SC_None);
+    KernelObjRef->setInit(MemberExprBases.back());
+    Stmt *DS = new (SemaRef.Context) DeclStmt(
+        DeclGroupRef(KernelObjRef), KernelCallerSrcLoc, KernelCallerSrcLoc);
+    BodyStmts.push_back(DS);
 
     // Replace references to the kernel object in kernel body, to use the
     // compiler generated local clone
     Stmt *NewBody =
-        replaceWithLocalClone(KernelCallerFunc->getParamDecl(0), KernelObjClone,
+        replaceWithLocalClone(KernelCallerFunc->getParamDecl(0), KernelObjRef,
                               KernelCallerFunc->getBody());
 
     // If kernel_handler argument is passed by SYCL kernel, replace references
@@ -2967,76 +2978,120 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
     return buildMemberExpr(DRE, ArrayField);
   }
 
-  // Creates an initialized entity for a field/item. In the case where this is a
-  // field, returns a normal member initializer, if we're in a sub-array of a MD
-  // array, returns an element initializer.
-  InitializedEntity getFieldEntity(FieldDecl *FD, QualType Ty) {
-    if (isArrayElement(FD, Ty))
-      return InitializedEntity::InitializeElement(SemaRef.getASTContext(),
-                                                  ArrayInfos.back().second,
-                                                  ArrayInfos.back().first);
-    return InitializedEntity::InitializeMember(FD, &VarEntity);
-  }
-
-  void addFieldInit(FieldDecl *FD, QualType Ty, MultiExprArg ParamRef) {
-    InitializationKind InitKind =
-        InitializationKind::CreateCopy(KernelCallerSrcLoc, KernelCallerSrcLoc);
-    addFieldInit(FD, Ty, ParamRef, InitKind);
-  }
-
-  void addFieldInit(FieldDecl *FD, QualType Ty, MultiExprArg ParamRef,
-                    InitializationKind InitKind) {
-    addFieldInit(FD, Ty, ParamRef, InitKind, getFieldEntity(FD, Ty));
+  // Build a placement new instruction to force the call to the default
+  // constructor if it exists. Sema has a BuildCXXNew method but we can't use it
+  // here as might be declared private.
+  void callFieldCtor() {
+    Expr *FieldExpr = MemberExprBases.back();
+    QualType ToType = FieldExpr->getType();
+    QualType CanonTy = ToType.getCanonicalType();
+    // This is gross but fine for clang. This should be moved to CodeGen anyway.
+    ToType.removeLocalConst();
+    ToType = SemaRef.Context.getPointerType(ToType);
+    FieldExpr = UnaryOperator::Create(
+        SemaRef.Context, FieldExpr, UO_AddrOf, ToType, VK_PRValue, OK_Ordinary,
+        KernelCallerSrcLoc, false, SemaRef.CurFPFeatureOverrides());
+
+    ExprResult NewExpr = SemaRef.BuildCXXNew(
+        KernelCallerSrcLoc, /*UseGlobal=*/true, KernelCallerSrcLoc, FieldExpr,
+        KernelCallerSrcLoc, {KernelCallerSrcLoc, KernelCallerSrcLoc}, CanonTy,
+        SemaRef.Context.getTrivialTypeSourceInfo(CanonTy, SourceLocation()),
+        std::nullopt /*FIXME: Array of accessors ?*/, {}, nullptr);
+
+    assert(!NewExpr.isInvalid() && "Can't build placement new!");
+    BodyStmts.push_back(NewExpr.getAs<Stmt>());
+  }
+
+  void doScalarInit(Expr* FromExpr, QualType Ty) {
+    assert(FromExpr->getType()->isScalarType());
+    // Compute the size of the memory buffer to be copied.
+    QualType SizeType = SemaRef.Context.getSizeType();
+    llvm::APInt Size(SemaRef.Context.getTypeSize(SizeType),
+                     SemaRef.Context.getTypeSizeInChars(Ty).getQuantity());
+
+    Expr *To = MemberExprBases.back();
+    QualType ToType = To->getType();
+    assert(ToType->isScalarType());
+    assert(FromExpr->getType().getTypePtr() == ToType.getTypePtr());
+
+    // This is gross but fine for clang.
+    // Could make it less gross by adding a const_cast but this should be moved
+    // to CodeGen anyway.
+    ToType.removeLocalConst();
+    ToType = SemaRef.Context.getPointerType(ToType);
+    To = BinaryOperator::Create(SemaRef.Context, To, FromExpr, BO_Assign, ToType,
+                                VK_LValue, OK_Ordinary, KernelCallerSrcLoc,
+                                SemaRef.CurFPFeatureOverrides());
+    BodyStmts.push_back(To);
+  }
+
+  void doMemCopyInit(Expr* FromExpr, QualType Ty) {
+    // Compute the size of the memory buffer to be copied.
+    QualType SizeType = SemaRef.Context.getSizeType();
+    llvm::APInt Size(SemaRef.Context.getTypeSize(SizeType),
+                     SemaRef.Context.getTypeSizeInChars(Ty).getQuantity());
+
+    Expr *From = UnaryOperator::Create(
+        SemaRef.Context, FromExpr, UO_AddrOf,
+        SemaRef.Context.getPointerType(FromExpr->getType()), VK_PRValue,
+        OK_Ordinary, KernelCallerSrcLoc, false,
+        SemaRef.CurFPFeatureOverrides());
+    Expr *To = MemberExprBases.back();
+    QualType ToType = To->getType();
+    // This is gross but fine for clang. This should be moved to CodeGen anyway.
+    ToType.removeLocalConst();
+    ToType = SemaRef.Context.getPointerType(ToType);
+    To = UnaryOperator::Create(SemaRef.Context, To, UO_AddrOf, ToType,
+                               VK_PRValue, OK_Ordinary, KernelCallerSrcLoc,
+                               false, SemaRef.CurFPFeatureOverrides());
+
+    Expr *CallArgs[] = {To, From,
+                        IntegerLiteral::Create(SemaRef.Context, Size, SizeType,
+                                               KernelCallerSrcLoc)};
+
+    ExprResult Call = SemaRef.BuildBuiltinCallExpr(
+        KernelCallerSrcLoc, Builtin::BI__builtin_memcpy, CallArgs);
+
+    assert(!Call.isInvalid() && "Call to __builtin_memcpy cannot fail!");
+    BodyStmts.push_back(Call.getAs<Stmt>());
+  }
+
+  void doInit(Expr* FromExpr, QualType Ty) {
+    if (FromExpr->getType()->isScalarType()) {
+      doScalarInit(FromExpr, Ty);
+    }
+    else {
+      doMemCopyInit(FromExpr, 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);
-
-    InitListExpr *ParentILE = CollectionInitExprs.back();
-    ParentILE->updateInit(SemaRef.getASTContext(), ParentILE->getNumInits(),
-                          Init.get());
+  Expr *getDeriveToBaseExpr(const CXXRecordDecl *Base,
+                            const CXXBaseSpecifier &BS, QualType Ty) {
+    CXXCastPath BasePath;
+    QualType DerivedTy(Base->getTypeForDecl(), 0);
+    QualType BaseTy = BS.getType();
+    SemaRef.CheckDerivedToBaseConversion(DerivedTy, BaseTy, KernelCallerSrcLoc,
+                                         SourceRange(), &BasePath,
+                                         /*IgnoreBaseAccess*/ true);
+    return ImplicitCastExpr::Create(
+        SemaRef.Context, BaseTy, CK_DerivedToBase, MemberExprBases.back(),
+        /* CXXCastPath=*/&BasePath, VK_LValue, FPOptionsOverride());
   }
 
-  void addBaseInit(const CXXBaseSpecifier &BS, QualType Ty,
-                   InitializationKind InitKind) {
-    InitializedEntity Entity = InitializedEntity::InitializeBase(
-        SemaRef.Context, &BS, /*IsInheritedVirtualBase*/ false, &VarEntity);
-    InitializationSequence InitSeq(SemaRef, Entity, InitKind, std::nullopt);
-    ExprResult Init = InitSeq.Perform(SemaRef, Entity, InitKind, std::nullopt);
+  void addSimpleBaseInit(const CXXRecordDecl *Base, const CXXBaseSpecifier &BS,
+                         QualType Ty) {
+    MemberExprBases.push_back(getDeriveToBaseExpr(Base, BS, Ty));
 
-    InitListExpr *ParentILE = CollectionInitExprs.back();
-    ParentILE->updateInit(SemaRef.getASTContext(), ParentILE->getNumInits(),
-                          Init.get());
-  }
-
-  void addBaseInit(const CXXBaseSpecifier &BS, QualType Ty,
-                   InitializationKind InitKind, MultiExprArg Args) {
-    InitializedEntity Entity = InitializedEntity::InitializeBase(
-        SemaRef.Context, &BS, /*IsInheritedVirtualBase*/ false, &VarEntity);
-    InitializationSequence InitSeq(SemaRef, Entity, InitKind, Args);
-    ExprResult Init = InitSeq.Perform(SemaRef, Entity, InitKind, Args);
+    Expr *ParamRef = createParamReferenceExpr();
 
-    InitListExpr *ParentILE = CollectionInitExprs.back();
-    ParentILE->updateInit(SemaRef.getASTContext(), ParentILE->getNumInits(),
-                          Init.get());
+    doMemCopyInit(ParamRef, Ty);
+    MemberExprBases.pop_back();
   }
 
-  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());
+  void addFieldInit(FieldDecl *FD, QualType Ty, Expr* ParamRef) {
+    addFieldMemberExpr(FD, Ty);
+    doInit(ParamRef, Ty);
+    removeFieldMemberExpr(FD, Ty);
   }
 
   // Adds an initializer that handles a simple initialization of a field.
@@ -3066,34 +3121,6 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
         SourceLocation(), SourceLocation(), SourceRange());
   }
 
-  void handleGeneratedType(FieldDecl *FD, QualType Ty) {
-    // Equivalent of the following code is generated here:
-    // void ocl_kernel(__generated_type GT) {
-    //   Kernel KernelObjClone { *(reinterpret_cast<UsersType*>(&GT)) };
-    // }
-
-    Expr *RCE = createReinterpretCastExpr(
-        createGetAddressOf(createParamReferenceExpr()),
-        SemaRef.Context.getPointerType(Ty));
-    Expr *Initializer = createDerefOp(RCE);
-    addFieldInit(FD, Ty, Initializer);
-  }
-
-  void handleGeneratedType(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS,
-                           QualType Ty) {
-    // Equivalent of the following code is generated here:
-    // void ocl_kernel(__generated_type GT) {
-    //   Kernel KernelObjClone { *(reinterpret_cast<UsersType*>(&GT)) };
-    // }
-    Expr *RCE = createReinterpretCastExpr(
-        createGetAddressOf(createParamReferenceExpr()),
-        SemaRef.Context.getPointerType(Ty));
-    Expr *Initializer = createDerefOp(RCE);
-    InitializationKind InitKind =
-        InitializationKind::CreateCopy(KernelCallerSrcLoc, KernelCallerSrcLoc);
-    addBaseInit(BS, Ty, InitKind, Initializer);
-  }
-
   MemberExpr *buildMemberExpr(Expr *Base, ValueDecl *Member) {
     DeclAccessPair MemberDAP = DeclAccessPair::make(Member, AS_none);
     MemberExpr *Result = SemaRef.BuildMemberExpr(
@@ -3147,41 +3174,78 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
         FPOptionsOverride()));
   }
 
-  // Creates an empty InitListExpr of the correct number of child-inits
-  // of this to append into.
-  void addCollectionInitListExpr(const CXXRecordDecl *RD) {
-    const ASTRecordLayout &Info =
-        SemaRef.getASTContext().getASTRecordLayout(RD);
-    uint64_t NumInitExprs = Info.getFieldCount() + RD->getNumBases();
-    addCollectionInitListExpr(QualType(RD->getTypeForDecl(), 0), NumInitExprs);
-  }
+  static void setupSpecialMemberType(ASTContext &Ctx, CXXMethodDecl *SpecialMem,
+                                     QualType ResultTy,
+                                     ArrayRef<QualType> Args) {
+    FunctionProtoType::ExtProtoInfo EPI;
 
-  InitListExpr *createInitListExpr(const CXXRecordDecl *RD) {
-    const ASTRecordLayout &Info =
-        SemaRef.getASTContext().getASTRecordLayout(RD);
-    uint64_t NumInitExprs = Info.getFieldCount() + RD->getNumBases();
-    return createInitListExpr(QualType(RD->getTypeForDecl(), 0), NumInitExprs);
+    EPI.ExceptionSpec.Type = EST_Unevaluated;
+    EPI.ExceptionSpec.SourceDecl = SpecialMem;
+
+    // Set the calling convention to the default for C++ instance methods.
+    EPI.ExtInfo = EPI.ExtInfo.withCallingConv(
+        Ctx.getDefaultCallingConvention(/*IsVariadic=*/false,
+                                        /*IsCXXMethod=*/true));
+
+    auto QT = Ctx.getFunctionType(ResultTy, Args, EPI);
+    SpecialMem->setType(QT);
   }
 
-  InitListExpr *createInitListExpr(QualType InitTy, uint64_t NumChildInits) {
-    InitListExpr *ILE = new (SemaRef.getASTContext()) InitListExpr(
-        SemaRef.getASTContext(), KernelCallerSrcLoc, {}, KernelCallerSrcLoc);
-    ILE->reserveInits(SemaRef.getASTContext(), NumChildInits);
-    ILE->setType(InitTy);
+  // Build an anonymous union class around the kernel object.
+  static CXXRecordDecl *getWrappingUnion(Sema &SemaRef, QualType KernelObj) {
+    ASTContext &Ctx = SemaRef.Context;
 
-    return ILE;
+    CanQualType CanClassType = Ctx.getCanonicalType(KernelObj);
+
+    CXXRecordDecl *WrapperClass = cast<CXXRecordDecl>(
+        Ctx.buildImplicitRecord("__wrapper_union", TTK_Union));
+    WrapperClass->startDefinition();
+    FieldDecl *KernelField = FieldDecl::Create(
+        Ctx, WrapperClass, SourceLocation(), SourceLocation(), /*Id=*/nullptr,
+        KernelObj, Ctx.getTrivialTypeSourceInfo(KernelObj, SourceLocation()),
+        /*BW=*/nullptr, /*Mutable=*/false, /*InitStyle=*/ICIS_NoInit);
+    KernelField->setAccess(AS_public);
+    WrapperClass->addDecl(KernelField);
+    // Build an empty DTor
+    {
+      DeclarationName Name =
+          Ctx.DeclarationNames.getCXXDestructorName(CanClassType);
+      DeclarationNameInfo NameInfo(Name, SourceLocation());
+      CXXDestructorDecl *DTor = CXXDestructorDecl::Create(
+          Ctx, WrapperClass, SourceLocation(), NameInfo, /*Type*/ QualType(),
+          /*TInfo=*/nullptr, /*isFPConstrained=*/false,
+          /*isInline=*/true, /*isImplicitlyDeclared=*/false,
+          ConstexprSpecKind::Constexpr);
+      DTor->setAccess(AS_public);
+      DTor->setTrivial(true);
+      setupSpecialMemberType(Ctx, DTor, Ctx.VoidTy, std::nullopt);
+      // Make an empty body
+      DTor->setBody(CompoundStmt::Create(Ctx, {}, FPOptionsOverride(), {}, {}));
+      WrapperClass->addDecl(DTor);
+    }
+    WrapperClass->completeDefinition();
+
+    return WrapperClass;
   }
 
-  // Create an empty InitListExpr of the type/size for the rest of the visitor
-  // to append into.
-  void addCollectionInitListExpr(QualType InitTy, uint64_t NumChildInits) {
+  static CXXRecordDecl *
+  createInUnionKernelWrapper(Sema &S, DeclContext *,
+                             const CXXRecordDecl *KernelObj) {
+    return getWrappingUnion(S, QualType(KernelObj->getTypeForDecl(), 0));
+  }
 
-    InitListExpr *ILE = createInitListExpr(InitTy, NumChildInits);
-    InitListExpr *ParentILE = CollectionInitExprs.back();
-    ParentILE->updateInit(SemaRef.getASTContext(), ParentILE->getNumInits(),
-                          ILE);
+  static VarDecl *
+  createInUnionKernelObjClone(Sema &S, DeclContext *DC,
+                              const CXXRecordDecl *WrappingUnion) {
+    ASTContext &Ctx = S.Context;
 
-    CollectionInitExprs.push_back(ILE);
+    VarDecl *VD = VarDecl::Create(Ctx, DC, SourceLocation(), SourceLocation(),
+                                  WrappingUnion->getIdentifier(),
+                                  QualType(WrappingUnion->getTypeForDecl(), 0),
+                                  nullptr, SC_None);
+    VD->setIsUsed();
+
+    return VD;
   }
 
   static VarDecl *createKernelObjClone(ASTContext &Ctx, DeclContext *DC,
@@ -3205,11 +3269,10 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
 
   // Default inits the type, then calls the init-method in the body.
   bool handleSpecialType(FieldDecl *FD, QualType Ty) {
-    addFieldInit(FD, Ty, std::nullopt,
-                 InitializationKind::CreateDefault(KernelCallerSrcLoc));
-
     addFieldMemberExpr(FD, Ty);
 
+    callFieldCtor();
+
     const auto *RecordDecl = Ty->getAsCXXRecordDecl();
     createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts);
     CXXMethodDecl *FinalizeMethod =
@@ -3223,10 +3286,16 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
     return true;
   }
 
-  bool handleSpecialType(const CXXBaseSpecifier &BS, QualType Ty) {
+  bool handleSpecialType(const CXXRecordDecl *Base,
+                         const CXXBaseSpecifier &BS, QualType Ty) {
+    MemberExprBases.push_back(getDeriveToBaseExpr(Base, BS, Ty));
+    callFieldCtor();
+
     const auto *RecordDecl = Ty->getAsCXXRecordDecl();
-    addBaseInit(BS, Ty, InitializationKind::CreateDefault(KernelCallerSrcLoc));
     createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts);
+
+    MemberExprBases.pop_back();
+
     return true;
   }
 
@@ -3276,92 +3345,37 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
     return IndexExpr.get();
   }
 
-  void addSimpleArrayInit(FieldDecl *FD, QualType FieldTy) {
-    Expr *ArrayRef = createSimpleArrayParamReferenceExpr(FieldTy);
-    InitializationKind InitKind = InitializationKind::CreateDirect({}, {}, {});
-
-    InitializedEntity Entity =
-        InitializedEntity::InitializeMember(FD, &VarEntity, /*Implicit*/ true);
-
-    addFieldInit(FD, FieldTy, ArrayRef, InitKind, Entity);
-  }
-
-  void addArrayElementInit(FieldDecl *FD, QualType T) {
-    Expr *RCE = createReinterpretCastExpr(
-        createGetAddressOf(ArrayParamBases.pop_back_val()),
-        SemaRef.Context.getPointerType(T));
-    Expr *Initializer = createDerefOp(RCE);
-    addFieldInit(FD, T, Initializer);
-  }
-
-  // This function is recursive in order to handle
-  // multi-dimensional arrays. If the array element is
-  // an array, it implies that the array is multi-dimensional.
-  // We continue recursion till we reach a non-array element to
-  // generate required array subscript expressions.
-  void createArrayInit(FieldDecl *FD, QualType T) {
-    const ConstantArrayType *CAT =
-        SemaRef.getASTContext().getAsConstantArrayType(T);
-
-    if (!CAT) {
-      addArrayElementInit(FD, T);
-      return;
-    }
-
-    QualType ET = CAT->getElementType();
-    uint64_t ElemCount = CAT->getSize().getZExtValue();
-    enterArray(FD, T, ET);
-
-    for (uint64_t Index = 0; Index < ElemCount; ++Index) {
-      ArrayInfos.back().second = Index;
-      Expr *ArraySubscriptExpr =
-          createArraySubscriptExpr(Index, ArrayParamBases.back());
-      ArrayParamBases.push_back(ArraySubscriptExpr);
-      createArrayInit(FD, ET);
-    }
-
-    leaveArray(FD, T, ET);
-  }
-
-  // This function is used to create initializers for a top
-  // level array which contains pointers. The openCl kernel
-  // parameter for this array will be a wrapper class
-  // which contains the generated type. This function generates
-  // code equivalent to:
-  // void ocl_kernel(__wrapper_class WrappedGT) {
-  //   Kernel KernelObjClone {
-  //   *reinterpret_cast<UserArrayET*>(&WrappedGT.GeneratedArr[0]),
-  //                           *reinterpret_cast<UserArrayET*>(&WrappedGT.GeneratedArr[1]),
-  //                           *reinterpret_cast<UserArrayET*>(&WrappedGT.GeneratedArr[2])
-  //                         };
-  // }
-  void handleGeneratedArrayType(FieldDecl *FD, QualType FieldTy) {
-    ArrayParamBases.push_back(createSimpleArrayParamReferenceExpr(FieldTy));
-    createArrayInit(FD, FieldTy);
-  }
-
 public:
   static constexpr const bool VisitInsideSimpleContainers = false;
   SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC,
                         const CXXRecordDecl *KernelObj,
                         FunctionDecl *KernelCallerFunc)
       : SyclKernelFieldHandler(S), DeclCreator(DC),
-        KernelObjClone(createKernelObjClone(S.getASTContext(),
-                                            DC.getKernelDecl(), KernelObj)),
+        WrappingUnion(
+            createInUnionKernelWrapper(S, DC.getKernelDecl(), KernelObj)),
+        // KernelObjClone(createKernelObjClone(S.getASTContext(),
+        //                                     DC.getKernelDecl(), KernelObj)),
+        // InUnionKernelObjClone(
+        //     createInUnionKernelObjClone(S, DC.getKernelDecl(), KernelObj)),
+        KernelObjClone(
+            createInUnionKernelObjClone(S, DC.getKernelDecl(), WrappingUnion)),
         VarEntity(InitializedEntity::InitializeVariable(KernelObjClone)),
         KernelObj(KernelObj), KernelCallerFunc(KernelCallerFunc),
         KernelCallerSrcLoc(KernelCallerFunc->getLocation()) {
-    CollectionInitExprs.push_back(createInitListExpr(KernelObj));
+    FieldDecl *WrappedField = *WrappingUnion->field_begin();
     annotateHierarchicalParallelismAPICalls();
 
     Stmt *DS = new (S.Context) DeclStmt(DeclGroupRef(KernelObjClone),
                                         KernelCallerSrcLoc, KernelCallerSrcLoc);
     BodyStmts.push_back(DS);
+
     DeclRefExpr *KernelObjCloneRef = DeclRefExpr::Create(
         S.Context, NestedNameSpecifierLoc(), KernelCallerSrcLoc, KernelObjClone,
-        false, DeclarationNameInfo(), QualType(KernelObj->getTypeForDecl(), 0),
-        VK_LValue);
+        false, DeclarationNameInfo(),
+        QualType(WrappingUnion->getTypeForDecl(), 0), VK_LValue);
+
     MemberExprBases.push_back(KernelObjCloneRef);
+    MemberExprBases.push_back(buildMemberExpr(KernelObjCloneRef, WrappedField));
   }
 
   ~SyclKernelBodyCreator() {
@@ -3373,9 +3387,9 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
     return handleSpecialType(FD, Ty);
   }
 
-  bool handleSyclSpecialType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
+  bool handleSyclSpecialType(const CXXRecordDecl *Base, const CXXBaseSpecifier &BS,
                              QualType Ty) final {
-    return handleSpecialType(BS, Ty);
+    return handleSpecialType(Base, BS, Ty);
   }
 
   bool handleSyclSpecConstantType(FieldDecl *FD, QualType Ty) final {
@@ -3390,32 +3404,21 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
   }
 
   bool handleSimpleArrayType(FieldDecl *FD, QualType FieldTy) final {
-    if (FD->hasAttr<SYCLGenerateNewTypeAttr>())
-      handleGeneratedArrayType(FD, FieldTy);
-    else
-      addSimpleArrayInit(FD, FieldTy);
+    Expr *ArrayRef = createSimpleArrayParamReferenceExpr(FieldTy);
+
+    addFieldInit(FD, FieldTy, ArrayRef);
     return true;
   }
 
   bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *FD,
                              QualType Ty) final {
-    CXXRecordDecl *RD = Ty->getAsCXXRecordDecl();
-    assert(RD && "Type must be a C++ record type");
-    if (RD->hasAttr<SYCLGenerateNewTypeAttr>())
-      handleGeneratedType(FD, Ty);
-    else
-      addSimpleFieldInit(FD, Ty);
+    addSimpleFieldInit(FD, Ty);
     return true;
   }
 
   bool handleNonDecompStruct(const CXXRecordDecl *RD,
                              const CXXBaseSpecifier &BS, QualType Ty) final {
-    CXXRecordDecl *BaseDecl = Ty->getAsCXXRecordDecl();
-    assert(BaseDecl && "Type must be a C++ record type");
-    if (BaseDecl->hasAttr<SYCLGenerateNewTypeAttr>())
-      handleGeneratedType(RD, BS, Ty);
-    else
-      addSimpleBaseInit(BS, Ty);
+    addSimpleBaseInit(RD, BS, Ty);
     return true;
   }
 
@@ -3453,7 +3456,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
 
   bool enterStruct(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final {
     ++StructDepth;
-    addCollectionInitListExpr(Ty->getAsCXXRecordDecl());
+    //addCollectionInitListExpr(Ty->getAsCXXRecordDecl());
 
     addFieldMemberExpr(FD, Ty);
     return true;
@@ -3461,7 +3464,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
 
   bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final {
     --StructDepth;
-    CollectionInitExprs.pop_back();
+    //CollectionInitExprs.pop_back();
 
     removeFieldMemberExpr(FD, Ty);
     return true;
@@ -3481,7 +3484,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
         SemaRef.Context, BaseTy, CK_DerivedToBase, MemberExprBases.back(),
         /* CXXCastPath=*/&BasePath, VK_LValue, FPOptionsOverride());
     MemberExprBases.push_back(Cast);
-    addCollectionInitListExpr(BaseTy->getAsCXXRecordDecl());
+
+    // addCollectionInitListExpr(BaseTy->getAsCXXRecordDecl());
     return true;
   }
 
@@ -3489,18 +3493,18 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
                    QualType) final {
     --StructDepth;
     MemberExprBases.pop_back();
-    CollectionInitExprs.pop_back();
+    //CollectionInitExprs.pop_back();
     return true;
   }
 
   bool enterArray(FieldDecl *FD, QualType ArrayType,
                   QualType ElementType) final {
-    const ConstantArrayType *CAT =
-        SemaRef.getASTContext().getAsConstantArrayType(ArrayType);
-    assert(CAT && "Should only be called on constant-size array.");
-    uint64_t ArraySize = CAT->getSize().getZExtValue();
-    addCollectionInitListExpr(ArrayType, ArraySize);
-    ArrayInfos.emplace_back(getFieldEntity(FD, ArrayType), 0);
+    // const ConstantArrayType *CAT =
+    //     SemaRef.getASTContext().getAsConstantArrayType(ArrayType);
+    // assert(CAT && "Should only be called on constant-size array.");
+    // uint64_t ArraySize = CAT->getSize().getZExtValue();
+    // addCollectionInitListExpr(ArrayType, ArraySize);
+    // ArrayInfos.emplace_back(getFieldEntity(FD, ArrayType), 0);
 
     // If this is the top-level array, we need to make a MemberExpr in addition
     // to an array subscript.
@@ -3509,7 +3513,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
   }
 
   bool nextElement(QualType, uint64_t Index) final {
-    ArrayInfos.back().second = Index;
+    //ArrayInfos.back().second = Index;
 
     // Pop off the last member expr base.
     if (Index != 0)
@@ -3522,8 +3526,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
 
   bool leaveArray(FieldDecl *FD, QualType ArrayType,
                   QualType ElementType) final {
-    CollectionInitExprs.pop_back();
-    ArrayInfos.pop_back();
+    //CollectionInitExprs.pop_back();
+    //ArrayInfos.pop_back();
 
     // Remove the IndexExpr.
     if (!FD->hasAttr<SYCLGenerateNewTypeAttr>())
diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
index 0467659cd5492..7c0293c3c13fa 100644
--- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -1,5 +1,8 @@
 #pragma once
 
+void* operator new  (__SIZE_TYPE__ size, void* ptr) noexcept;
+void* operator new[](__SIZE_TYPE__ size, void* ptr) noexcept;
+
 #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
 #define __SYCL_TYPE(x) [[__sycl_detail__::sycl_type(x)]]
 
diff --git a/clang/test/CodeGenSYCL/accessor-readonly-invalid-lib.cpp b/clang/test/CodeGenSYCL/accessor-readonly-invalid-lib.cpp
index 7ca8bbb86a205..e1ea376073159 100644
--- a/clang/test/CodeGenSYCL/accessor-readonly-invalid-lib.cpp
+++ b/clang/test/CodeGenSYCL/accessor-readonly-invalid-lib.cpp
@@ -2,6 +2,8 @@
 //
 // Test which verifies that readonly attribute is generated for unexpected access mode value.
 
+void* operator new  (__SIZE_TYPE__ size, void* ptr) noexcept;
+
 // Dummy library with unexpected access::mode enum value.
 namespace sycl {
 inline namespace _V1 {
diff --git a/clang/test/CodeGenSYCL/accessor_inheritance.cpp b/clang/test/CodeGenSYCL/accessor_inheritance.cpp
index 732ee2cb5e6a7..7e75faad7275f 100644
--- a/clang/test/CodeGenSYCL/accessor_inheritance.cpp
+++ b/clang/test/CodeGenSYCL/accessor_inheritance.cpp
@@ -50,7 +50,11 @@ int main() {
 // CHECK: [[ARG_C]].addr.ascast = addrspacecast ptr [[ARG_C]].addr to ptr addrspace(4)
 //
 // Lambda object alloca
-// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_.]+]] = addrspacecast ptr [[KERNEL]] to ptr addrspace(4)
+// CHECK: [[KERNEL:%[a-zA-Z0-9_]+]] = alloca %union.__wrapper_union
+// CHECK: [[KERNEL_UNION_OBJ:%[a-zA-Z0-9_.]+]] = addrspacecast %union.__wrapper_union* [[KERNEL]] to %union.__wrapper_union addrspace(4)*
+// CHECK: [[KERNEL_OBJ_PTR_ALLOCA:%[a-zA-Z0-9_.]+]] = alloca %class{{.*}}.anon addrspace(4)*, align 8
+// CHECK: [[KERNEL_OBJ_PTR:%[a-zA-Z0-9_.]+]] = addrspacecast %class.anon addrspace(4)** [[KERNEL_OBJ_PTR_ALLOCA]] to %class.anon addrspace(4)* addrspace(4)*
+
 //
 // Kernel argument stores
 // CHECK: store i32 [[ARG_A]], ptr addrspace(4) [[ARG_A]].addr.ascast
@@ -60,13 +64,15 @@ int main() {
 // CHECK: store i32 [[ARG_C]], ptr addrspace(4) [[ARG_C]].addr.ascast
 //
 // Check A and B scalar fields initialization
-// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, ptr addrspace(4) [[KERNEL_OBJ]], i32 0, i32 0
-// CHECK: [[FIELD_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, ptr addrspace(4) [[GEP]], i32 0, i32 0
-// CHECK: [[ARG_A_LOAD:%[a-zA-Z0-9_]+]] = load i32, ptr addrspace(4) [[ARG_A]].addr.ascast
-// CHECK: store i32 [[ARG_A_LOAD]], ptr addrspace(4) [[FIELD_A]]
-// CHECK: [[FIELD_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, ptr addrspace(4) [[GEP]], i32 0, i32 1
-// CHECK: [[ARG_B_LOAD:%[a-zA-Z0-9_]+]] = load i32, ptr addrspace(4) [[ARG_B]].addr.ascast
-// CHECK: store i32 [[ARG_B_LOAD]], ptr addrspace(4) [[FIELD_B]]
+// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_.]+]] = bitcast %union.__wrapper_union addrspace(4)* [[KERNEL_UNION_OBJ]] to %class{{.*}}.anon addrspace(4)*
+// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
+// CHECK: [[BITCAST:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP]] to %struct{{.*}}Base addrspace(4)*
+// CHECK: [[FIELD_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST]], i32 0, i32 0
+// CHECK: [[ARG_A_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32 addrspace(4)* [[ARG_A]].addr.ascast
+// CHECK: store i32 [[ARG_A_LOAD]], i32 addrspace(4)* [[FIELD_A]]
+// CHECK: [[FIELD_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST]], i32 0, i32 1
+// CHECK: [[ARG_B_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32 addrspace(4)* [[ARG_B]].addr.ascast
+// CHECK: store i32 [[ARG_B_LOAD]], i32 addrspace(4)* [[FIELD_B]]
 //
 // Check accessors initialization
 // CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, ptr addrspace(4) [[GEP]], i32 0, i32 2
@@ -82,11 +88,15 @@ int main() {
 // CHECK: store i32 [[ARG_C_LOAD]], ptr addrspace(4) [[FIELD_C]]
 //
 // Check __init method calls
-// CHECK: [[GEP2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, ptr addrspace(4) [[KERNEL_OBJ]], i32 0, i32 0
-// CHECK: [[ACC1_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, ptr addrspace(4) [[GEP2]], i32 0, i32 2
-// CHECK: [[ACC1_DATA_LOAD:%[a-zA-Z0-9_]+]] = load ptr addrspace(1), ptr addrspace(4) [[ACC1_DATA]].addr.ascast
-// CHECK: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{[^,]*}} [[ACC1_FIELD]], ptr addrspace(1) noundef [[ACC1_DATA_LOAD]]
+// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_]+]] = load %class.anon addrspace(4)*, %class.anon addrspace(4)* addrspace(4)* [[KERNEL_OBJ_PTR]]
+// CHECK: [[GEP2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
+// CHECK: [[BITCAST3:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP2]] to %struct{{.*}}Base addrspace(4)*
+// CHECK: [[ACC1_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST3]], i32 0, i32 2
+// CHECK: [[ACC1_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(4)* [[ACC1_DATA]].addr.ascast
+// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC1_FIELD]], i8 addrspace(1)* [[ACC1_DATA_LOAD]]
 //
-// CHECK: [[GEP3:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, ptr addrspace(4) [[KERNEL_OBJ]], i32 0, i32 0
-// CHECK: [[ACC2_DATA_LOAD:%[a-zA-Z0-9_]+]] = load ptr addrspace(1), ptr addrspace(4) [[ACC2_DATA]].addr.ascast
-// CHECK: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{[^,]*}}, ptr addrspace(1) noundef [[ACC2_DATA_LOAD]]
+// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_]+]] = load %class.anon addrspace(4)*, %class.anon addrspace(4)* addrspace(4)* [[KERNEL_OBJ_PTR]]
+// CHECK: [[GEP3:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
+// CHECK: [[ACC2_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(4)* [[ACC2_DATA]].addr.ascast
+// CHECK: [[BITCAST4:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP3]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
+// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[BITCAST4]], i8 addrspace(1)* [[ACC2_DATA_LOAD]]
diff --git a/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp b/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp
index ba1e7e4a10ff4..1fc7d6311c339 100644
--- a/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp
+++ b/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp
@@ -1,4 +1,8 @@
 // RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -opaque-pointers -emit-llvm %s -o - | FileCheck %s
+
+// Set the cursor so that CEHCK-DAG don't pickup wrong matches
+// CHECK: define {{.*}} @_Z6usagesv
+
 void bar(int & Data) {}
 // CHECK-DAG: define {{.*}}spir_func void @[[RAW_REF:[a-zA-Z0-9_]+]](ptr addrspace(4) noundef align 4 dereferenceable(4) %
 void bar2(int & Data) {}
diff --git a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
index 801af7de09965..0aa1b5cd18746 100644
--- a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
+++ b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
@@ -27,7 +27,10 @@ int main() {
 // Check alloca for pointer argument
 // CHECK: [[MEM_ARG]].addr = alloca ptr addrspace(1)
 // Check lambda object alloca
-// CHECK: [[ANONALLOCA:%[a-zA-Z0-9_]+]] = alloca %class.anon
+// CHECK: [[UNIONALLOCA:%[a-zA-Z0-9_]+]] = alloca %union.__wrapper_union
+// CHECK: [[UNION:%[a-zA-Z0-9_.]+]] = addrspacecast %union.__wrapper_union* [[UNIONALLOCA]] to %union.__wrapper_union addrspace(4)*
+// CHECK: [[ANONPTRALLOCA:%[a-zA-Z0-9_.]+]] = alloca %class{{.*}}.anon addrspace(4)*, align 8
+// CHECK: [[ANONPTRALLOCA_PTR:%[a-zA-Z0-9_.]+]] = addrspacecast %class.anon addrspace(4)** [[ANONPTRALLOCA]] to %class.anon addrspace(4)* addrspace(4)*
 // Check allocas for ranges
 // CHECK: [[ARANGEA:%agg.tmp.*]] = alloca %"struct.sycl::_V1::range"
 // CHECK: [[MRANGEA:%agg.tmp.*]] = alloca %"struct.sycl::_V1::range"
@@ -44,7 +47,8 @@ int main() {
 // CHECK: call spir_func {{.*}}accessor
 
 // Check accessor GEP
-// CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, ptr addrspace(4) [[ANON]], i32 0, i32 0
+// CHECK: [[ANON:%[0-9]+]] = load %class.anon addrspace(4)*, %class.anon addrspace(4)* addrspace(4)* [[ANONPTRALLOCA_PTR]]
+// CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* [[ANON]], i32 0, i32 0
 
 // Check load from kernel pointer argument alloca
 // CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load ptr addrspace(1), ptr addrspace(4) [[MEM_ARG]].addr.ascast
diff --git a/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp b/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp
index 5e3fca15827c9..b2fdbf003128a 100644
--- a/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp
+++ b/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp
@@ -317,7 +317,7 @@ int main() {
 
     // Test attribute is not propagated.
     // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name32() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
-    // CHECK: define {{.*}}spir_func void @{{.*}}Functor10{{.*}}(ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) %this) #3 comdat align 2
+    // CHECK: define {{.*}}spir_func void @{{.*}}Functor10{{.*}}(ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) %this) #2 comdat align 2
     // CHECK-NOT: noalias
     // CHECK-SAME: {
     // CHECK: define dso_local spir_func void @_Z4foo8v()
@@ -325,12 +325,12 @@ int main() {
     h.single_task<class kernel_name32>(f10);
 
     // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name33() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
-    // CHECK: define {{.*}}spir_func void @{{.*}}Foo8{{.*}}(ptr addrspace(4) noalias noundef align 1 dereferenceable_or_null(1) %this) #3 comdat align 2
+    // CHECK: define {{.*}}spir_func void @{{.*}}Foo8{{.*}}(ptr addrspace(4) noalias noundef align 1 dereferenceable_or_null(1) %this) #2 comdat align 2
     Foo8 boo8;
     h.single_task<class kernel_name33>(boo8);
 
     // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name34() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
-    // CHECK: define {{.*}}spir_func void @{{.*}}(ptr addrspace(4) noalias noundef align 1 dereferenceable_or_null(1) %this) #4 align 2
+    // CHECK: define {{.*}}spir_func void @{{.*}}(ptr addrspace(4) noalias noundef align 1 dereferenceable_or_null(1) %this) #3 align 2
     h.single_task<class kernel_name34>(
         []() [[intel::kernel_args_restrict]]{});
 
diff --git a/clang/test/CodeGenSYCL/generated-types-initialization.cpp b/clang/test/CodeGenSYCL/generated-types-initialization.cpp
index b6e3e160ac2dd..f0a30f8430f25 100644
--- a/clang/test/CodeGenSYCL/generated-types-initialization.cpp
+++ b/clang/test/CodeGenSYCL/generated-types-initialization.cpp
@@ -41,8 +41,10 @@ int main() {
 // CHECK: define dso_local spir_kernel void @{{.*}}basic(ptr noundef byval(%struct.__generated_B) align 8 %_arg_Obj)
 //
 // Kernel object clone.
-// CHECK: %[[K:[a-zA-Z0-9_.]+]] = alloca %class.anon
+// CHECK: %[[K:[a-zA-Z0-9_.]+]] = alloca %union.__wrapper_union
+// CHECK: %[[K_PTR_ALLOCA:[a-zA-Z0-9_.]+]] = alloca ptr addrspace(4)
 // CHECK: %[[K_as_cast:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[K]] to ptr addrspace(4)
+// CHECK: %[[K_PTR_ALLOCA_as_cast:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[K_PTR_ALLOCA]] to ptr addrspace(4)
 //
 // Argument reference.
 // CHECK: %[[Arg_ref:[a-zA-Z0-9_.]+]] = addrspacecast ptr %_arg_Obj to ptr addrspace(4)
@@ -52,20 +54,26 @@ int main() {
 // CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %[[GEP]], ptr addrspace(4) align 8 %[[Arg_ref]], i64 16, i1 false)
 //
 // Kernel body call.
-// CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(16) %[[K_as_cast]])
+// CHECK: store ptr addrspace(4) %[[K_as_cast]], ptr addrspace(4) %[[K_PTR_ALLOCA_as_cast]]
+// CHECK: %[[K_PTR:[a-zA-Z0-9_.]+]] = load ptr addrspace(4), ptr addrspace(4) %[[K_PTR_ALLOCA_as_cast]]
+// CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(16) %[[K_PTR]])
 
 // CHECK: define dso_local spir_kernel void @{{.*}}nns(ptr noundef byval(%struct.__generated_B.0) align 8 %_arg_NNSObj)
 //
 // Kernel object clone.
-// CHECK: %[[NNSK:[a-zA-Z0-9_.]+]] = alloca %class.anon.2
+// CHECK: %[[NNSK:[a-zA-Z0-9_.]+]] = alloca %union.__wrapper_union.2
+// CHECK: %[[NNSK_PTR_ALLOCA:[a-zA-Z0-9_.]+]] = alloca ptr addrspace(4)
 // CHECK: %[[NNSK_as_cast:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[NNSK]] to ptr addrspace(4)
+// CHECK: %[[NNSK_PTR_ALLOCA_as_cast:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[NNSK_PTR_ALLOCA]] to ptr addrspace(4)
 //
 // Argument reference.
 // CHECK: %[[NNSArg_ref:[a-zA-Z0-9_.]+]] = addrspacecast ptr %_arg_NNSObj to ptr addrspace(4)
 //
 // Initialization.
-// CHECK: %[[NNSGEP:[a-zA-Z0-9_.]+]] = getelementptr inbounds %class.anon.2, ptr addrspace(4) %[[NNSK_as_cast]], i32 0, i32 0
+// CHECK: %[[NNSGEP:[a-zA-Z0-9_.]+]] = getelementptr inbounds %class.anon.3, ptr addrspace(4) %[[NNSK_as_cast]], i32 0, i32 0
 // CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %[[NNSGEP]], ptr addrspace(4) align 8 %[[NNSArg_ref]], i64 16, i1 false)
 //
 // Kernel body call.
-// CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_ENKUlvE_clEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(16) %[[NNSK_as_cast]])
+// CHECK: store ptr addrspace(4) %[[NNSK_as_cast]], ptr addrspace(4) %[[NNSK_PTR_ALLOCA_as_cast]]
+// CHECK: %[[NNSK_PTR:[a-zA-Z0-9_.]+]] = load ptr addrspace(4), ptr addrspace(4) %[[NNSK_PTR_ALLOCA_as_cast]]
+// CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_ENKUlvE_clEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(16) %[[NNSK_PTR]])
diff --git a/clang/test/CodeGenSYCL/image_accessor.cpp b/clang/test/CodeGenSYCL/image_accessor.cpp
index c697407c72059..3c3fd11a01c01 100644
--- a/clang/test/CodeGenSYCL/image_accessor.cpp
+++ b/clang/test/CodeGenSYCL/image_accessor.cpp
@@ -7,22 +7,22 @@
 // RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-3DWO
 //
 // CHECK-1DRO: define {{.*}}spir_kernel void @{{.*}}(ptr addrspace(1) [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
-// CHECK-1DRO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z]+}}, ptr addrspace(1) %{{[0-9]+}})
+// CHECK-1DRO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z0-9]+}}, ptr addrspace(1) %{{[0-9]+}})
 //
 // CHECK-2DRO: define {{.*}}spir_kernel void @{{.*}}(ptr addrspace(1) [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
-// CHECK-2DRO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z]+}}, ptr addrspace(1) %{{[0-9]+}})
+// CHECK-2DRO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z0-9]+}}, ptr addrspace(1) %{{[0-9]+}})
 //
 // CHECK-3DRO: define {{.*}}spir_kernel void @{{.*}}(ptr addrspace(1) [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
-// CHECK-3DRO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z]+}}, ptr addrspace(1) %{{[0-9]+}})
+// CHECK-3DRO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z0-9]+}}, ptr addrspace(1) %{{[0-9]+}})
 //
 // CHECK-1DWO: define {{.*}}spir_kernel void @{{.*}}(ptr addrspace(1) [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
-// CHECK-1DWO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z]+}}, ptr addrspace(1) %{{[0-9]+}})
+// CHECK-1DWO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z0-9]+}}, ptr addrspace(1) %{{[0-9]+}})
 //
 // CHECK-2DWO: define {{.*}}spir_kernel void @{{.*}}(ptr addrspace(1) [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
-// CHECK-2DWO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z]+}}, ptr addrspace(1) %{{[0-9]+}})
+// CHECK-2DWO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z0-9]+}}, ptr addrspace(1) %{{[0-9]+}})
 //
 // CHECK-3DWO: define {{.*}}spir_kernel void @{{.*}}(ptr addrspace(1) [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
-// CHECK-3DWO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z]+}}, ptr addrspace(1) %{{[0-9]+}})
+// CHECK-3DWO: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{.*}} %{{[a-zA-Z0-9]+}}, ptr addrspace(1) %{{[0-9]+}})
 //
 // TODO: Add tests for the image_array opencl datatype support.
 #include "Inputs/sycl.hpp"
diff --git a/clang/test/CodeGenSYCL/inheritance.cpp b/clang/test/CodeGenSYCL/inheritance.cpp
index 2806779730e9a..da013c6190f6c 100644
--- a/clang/test/CodeGenSYCL/inheritance.cpp
+++ b/clang/test/CodeGenSYCL/inheritance.cpp
@@ -56,24 +56,24 @@ int main() {
 
 // Check allocas for kernel parameters and local functor object
 // CHECK: %[[ARG_A_ALLOCA:[a-zA-Z0-9_.]+]] = alloca i32, align 4
-// CHECK: %[[LOCAL_OBJECT_ALLOCA:[a-zA-Z0-9_.]+]] = alloca %struct.derived, align 8
+// CHECK: %[[UNIONALLOCA:[a-zA-Z0-9_]+]] = alloca %union.__wrapper_union
+// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = alloca ptr addrspace(4), align 8
 // CHECK: %[[ARG_A:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[ARG_A_ALLOCA]] to ptr addrspace(4)
-// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[LOCAL_OBJECT_ALLOCA]] to ptr addrspace(4)
+// CHECK: %[[UNION:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[UNIONALLOCA]] to ptr addrspace(4)
 // CHECK: %[[ARG_BASE:[a-zA-Z0-9_.]+]] = addrspacecast ptr %_arg__base to ptr addrspace(4)
 // CHECK: %[[ARG_BASE1:[a-zA-Z0-9_.]+]] = addrspacecast ptr %_arg__base1 to ptr addrspace(4)
 // CHECK: store i32 %_arg_a, ptr addrspace(4) %[[ARG_A]], align 4
 
 // Initialize 'base' subobject
-// CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %[[LOCAL_OBJECT]], ptr addrspace(4) align 4 %[[ARG_BASE]], i64 12, i1 false)
+// CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %[[UNION]], ptr addrspace(4) align 4 %[[ARG_BASE]], i64 12, i1 false)
 
 // Initialize 'second_base' subobject
 // First, derived-to-base cast with offset:
-// CHECK: %[[OFFSET_CALC:.*]] = getelementptr inbounds i8, ptr addrspace(4) %[[LOCAL_OBJECT]], i64 16
+// CHECK: %[[OFFSET_CALC:.*]] = getelementptr inbounds i8, ptr addrspace(4) %[[UNION]], i64 16
 // Initialize 'second_base'
 // CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %[[OFFSET_CALC]], ptr addrspace(4) align 8 %[[ARG_BASE1]], i64 24, i1 false)
 
 // Initialize field 'a'
-// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.derived, ptr addrspace(4) %[[LOCAL_OBJECT]], i32 0, i32 3
 // CHECK: %[[LOAD_A:[0-9]+]] = load i32, ptr addrspace(4) %[[ARG_A]], align 4
+// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.derived, ptr addrspace(4) %[[UNION]], i32 0, i32 3
 // CHECK: store i32 %[[LOAD_A]], ptr addrspace(4) %[[GEP_A]]
-
diff --git a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp
index ebe9e383e0e3e..2d2de8e561cd3 100644
--- a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp
@@ -39,7 +39,7 @@ int main() {
 // CHECK: [[MEM_ARG2:%[a-zA-Z0-9_.]+]] = alloca ptr addrspace(1), align 8
 
 // CHECK lambda object alloca
-// CHECK: [[LOCAL_OBJECTA:%__SYCLKernel]] = alloca %class.anon, align 4
+// CHECK: [[UNION_ALLOCA:%.*]] = alloca %union.__wrapper_union, align 4
 
 // CHECK allocas for ranges
 // CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::range"
@@ -48,9 +48,10 @@ int main() {
 // CHECK: [[ACC_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::range"
 // CHECK: [[MEM_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::range"
 // CHECK: [[OFFSET2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::id"
+// CHECK: [[FUNCTOR_PTRALLOCA:%.*]] = alloca ptr addrspace(4)
 
 // CHECK lambda object addrspacecast
-// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast ptr [[LOCAL_OBJECTA]] to ptr addrspace(4)
+// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast ptr [[UNION_ALLOCA]] to ptr addrspace(4)
 
 // CHECK addrspacecasts for ranges
 // CHECK: [[ACC_RANGE1AS:%.*]] = addrspacecast ptr [[ACC_RANGE1A]] to ptr addrspace(4)
@@ -59,16 +60,13 @@ int main() {
 // CHECK: [[ACC_RANGE2AS:%.*]] = addrspacecast ptr [[ACC_RANGE2A]] to ptr addrspace(4)
 // CHECK: [[MEM_RANGE2AS:%.*]] = addrspacecast ptr [[MEM_RANGE2A]] to ptr addrspace(4)
 // CHECK: [[OFFSET2AS:%.*]] = addrspacecast ptr [[OFFSET2A]] to ptr addrspace(4)
+
 // CHECK accessor array default inits
 // CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, ptr addrspace(4) [[LOCAL_OBJECT]], i32 0, i32 0
 // CHECK: [[BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], ptr addrspace(4) [[ACCESSOR_ARRAY1]], i64 0, i64 0
 // Clang takes advantage of element 1 having the same address as the array, so it doesn't do a GEP.
 // CTOR Call #1
 // CHECK: call spir_func void @{{.+}}(ptr addrspace(4) {{[^,]*}} [[BEGIN]])
-// CHECK: [[ELEM2_GEP:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [[ACCESSOR]], ptr addrspace(4) [[BEGIN]], i64 1
-// CTOR Call #2
-// CHECK: call spir_func void @{{.+}}(ptr addrspace(4) {{[^,]*}} [[ELEM2_GEP]])
-
 // CHECK acc[0] __init method call
 // CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, ptr addrspace(4) [[LOCAL_OBJECT]], i32 0, i32 0
 // CHECK: [[INDEX1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], ptr addrspace(4) [[ACCESSOR_ARRAY1]], i64 0, i64 0
@@ -79,6 +77,11 @@ int main() {
 // CHECK: [[OFFSET1:%.*]] = addrspacecast ptr addrspace(4) [[OFFSET1AS]] to ptr
 // CHECK: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{[^,]*}} [[INDEX1]], ptr addrspace(1) noundef [[MEM_LOAD1]], ptr noundef byval({{.*}}) align 4 [[ACC_RANGE1]], ptr noundef byval({{.*}}) align 4 [[MEM_RANGE1]], ptr noundef byval({{.*}}) align 4 [[OFFSET1]])
 
+// CTOR Call #2
+// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, ptr addrspace(4) [[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: [[ELEM2_GEP:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [2 x [[ACCESSOR]]], ptr addrspace(4) [[ACCESSOR_ARRAY2]], i64 0, i64 1
+// CHECK: call spir_func void @{{.+}}(ptr addrspace(4) {{[^,]*}} [[ELEM2_GEP]])
+
 // CHECK acc[1] __init method call
 // CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, ptr addrspace(4) [[LOCAL_OBJECT]], i32 0, i32 0
 // CHECK: [[INDEX2:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], ptr addrspace(4) [[ACCESSOR_ARRAY2]], i64 0, i64 1
diff --git a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp
index 13f1d28daba5d..6744264af6512 100644
--- a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp
@@ -42,7 +42,7 @@ int main() {
 // CHECK: [[MEM_ARG1]].addr{{[0-9]*}} = alloca ptr addrspace(1), align 8
 
 // Check lambda object alloca
-// CHECK: [[LOCAL_OBJECTA:%__SYCLKernel]] = alloca %class{{.*}}.anon, align 4
+// CHECK: [[LOCAL_OBJECTA:%__wrapper_union]] = alloca %union.__wrapper_union, align 4
 
 // Check allocas for ranges
 // CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::range"
@@ -53,7 +53,7 @@ int main() {
 // CHECK: [[OFFSET2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::id"
 
 // Check lambda object addrspacecast
-// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast ptr %__SYCLKernel to ptr addrspace(4)
+// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast ptr [[LOCAL_OBJECTA]] to ptr addrspace(4)
 
 // Check addrspacecast for ranges
 // CHECK: [[ACC_RANGE1AS:%.*]] = addrspacecast ptr [[ACC_RANGE1A]] to ptr addrspace(4)
diff --git a/clang/test/CodeGenSYCL/kernel_binding_decls.cpp b/clang/test/CodeGenSYCL/kernel_binding_decls.cpp
index 21354f9348804..32b2129cee1a2 100644
--- a/clang/test/CodeGenSYCL/kernel_binding_decls.cpp
+++ b/clang/test/CodeGenSYCL/kernel_binding_decls.cpp
@@ -25,22 +25,24 @@ void foo() {
 // Check alloca of the captured types
 // CHECK: %_arg_x.addr = alloca i32, align 4
 // CHECK: %_arg_f2.addr = alloca float, align 4
-// CHECK: %__SYCLKernel = alloca %class.anon, align 4
+// CHECK: %__wrapper_union = alloca %union.__wrapper_union, align 4
 
 // Copy the parameters into the alloca-ed addresses
 // CHECK: store i32 %_arg_x, ptr addrspace(4) %_arg_x.addr
 // CHECK: store float %_arg_f2, ptr addrspace(4) %_arg_f2.addr
 
 // Store the int and the float into the struct created
-// CHECK: %x = getelementptr inbounds %class.anon, ptr addrspace(4) %__SYCLKernel{{.*}}, i32 0, i32 0
-// CHECK: %0 = load i32, ptr addrspace(4) %_arg_x.addr
-// CHECK: store i32 %0, ptr addrspace(4) %x
-// CHECK: %f2 = getelementptr inbounds %class.anon, ptr addrspace(4) %__SYCLKernel{{.*}}, i32 0, i32 1
-// CHECK: %1 = load float, ptr addrspace(4) %_arg_f2.addr
-// CHECK: store float %1, ptr addrspace(4) %f2
+// CHECK: %[[X_VALUE:[A-Za-z0-9]*]] = load i32, ptr addrspace(4) %_arg_x.addr
+// CHECK: %x = getelementptr inbounds %class.anon, ptr addrspace(4) %__wrapper_union{{.*}}, i32 0, i32 0
+// CHECK: store i32 %[[X_VALUE]], ptr addrspace(4) %x
+// CHECK: %[[F2_VALUE:[A-Za-z0-9]*]] = load float, ptr addrspace(4) %_arg_f2.addr
+// CHECK: %f2 = getelementptr inbounds %class.anon, ptr addrspace(4) %__wrapper_union{{.*}}, i32 0, i32 1
+// CHECK: store float %[[F2_VALUE]], ptr addrspace(4) %f2
 
 // Call the lambda
-// CHECK: call spir_func void @{{.*}}foo{{.*}}(ptr addrspace(4) {{.*}} %__SYCLKernel{{.*}})
+// CHECK: store ptr addrspace(4) %__wrapper_union{{.*}}, ptr addrspace(4) %[[KERNEL_REF_ADDR:[A-Za-z0-9]*]]
+// CHECK: %[[KERNEL_REF:[A-Za-z0-9]*]] = load ptr addrspace(4), ptr addrspace(4) %[[KERNEL_REF_ADDR]]
+// CHECK: call spir_func void @{{.*}}foo{{.*}}(ptr addrspace(4) {{.*}} %[[KERNEL_REF]])
 // CHECK:   ret void
 
 // Check the lambda call
diff --git a/clang/test/CodeGenSYCL/no_opaque_accessor_inheritance.cpp b/clang/test/CodeGenSYCL/no_opaque_accessor_inheritance.cpp
index 96f36f11c727a..67720b2c3f059 100644
--- a/clang/test/CodeGenSYCL/no_opaque_accessor_inheritance.cpp
+++ b/clang/test/CodeGenSYCL/no_opaque_accessor_inheritance.cpp
@@ -42,7 +42,8 @@ int main() {
 // CHECK: [[ACC1_DATA]].addr = alloca i8 addrspace(1)
 // CHECK: [[ACC2_DATA]].addr = alloca i8 addrspace(1)*
 // CHECK: [[ARG_C]].addr = alloca i32
-// CHECK: [[KERNEL:%[a-zA-Z0-9_]+]] = alloca %class{{.*}}.anon
+// CHECK: [[KERNEL:%[a-zA-Z0-9_]+]] = alloca %union.__wrapper_union
+// CHECK: [[KERNEL_PTR:%[a-zA-Z0-9_]+]] = alloca %class{{.*}}.anon addrspace(4)*
 // CHECK: [[ARG_A]].addr.ascast = addrspacecast i32* [[ARG_A]].addr to i32 addrspace(4)*
 // CHECK: [[ARG_B]].addr.ascast = addrspacecast i32* [[ARG_B]].addr to i32 addrspace(4)*
 // CHECK: [[ACC1_DATA]].addr.ascast = addrspacecast i8 addrspace(1)** [[ACC1_DATA]].addr to i8 addrspace(1)* addrspace(4)*
@@ -50,7 +51,7 @@ int main() {
 // CHECK: [[ARG_C]].addr.ascast = addrspacecast i32* [[ARG_C]].addr to i32 addrspace(4)*
 //
 // Lambda object alloca
-// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_.]+]] = addrspacecast %class{{.*}}.anon* [[KERNEL]] to %class{{.*}}.anon addrspace(4)*
+// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_.]+]] = addrspacecast %union.__wrapper_union* [[KERNEL]] to %union.__wrapper_union addrspace(4)*
 //
 // Kernel argument stores
 // CHECK: store i32 [[ARG_A]], i32 addrspace(4)* [[ARG_A]].addr.ascast
@@ -60,24 +61,45 @@ int main() {
 // CHECK: store i32 [[ARG_C]], i32 addrspace(4)* [[ARG_C]].addr.ascast
 //
 // Check A and B scalar fields initialization
-// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
+// CHECK: [[ARG_A_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32 addrspace(4)* [[ARG_A]].addr.ascast
+// CHECK: [[KERNEL_BC:%[a-zA-Z0-9_]+]] = bitcast %union.__wrapper_union addrspace(4)* [[KERNEL_OBJ]] to %class.anon addrspace(4)*
+// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_BC]], i32 0, i32 0
 // CHECK: [[BITCAST:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP]] to %struct{{.*}}Base addrspace(4)*
 // CHECK: [[FIELD_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST]], i32 0, i32 0
-// CHECK: [[ARG_A_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32 addrspace(4)* [[ARG_A]].addr.ascast
 // CHECK: store i32 [[ARG_A_LOAD]], i32 addrspace(4)* [[FIELD_A]]
-// CHECK: [[FIELD_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST]], i32 0, i32 1
 // CHECK: [[ARG_B_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32 addrspace(4)* [[ARG_B]].addr.ascast
+// CHECK: [[KERNEL_BC:%[a-zA-Z0-9_]+]] = bitcast %union.__wrapper_union addrspace(4)* [[KERNEL_OBJ]] to %class.anon addrspace(4)*
+// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_BC]], i32 0, i32 0
+// CHECK: [[BITCAST:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP]] to %struct{{.*}}Base addrspace(4)*
+// CHECK: [[FIELD_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST]], i32 0, i32 1
 // CHECK: store i32 [[ARG_B_LOAD]], i32 addrspace(4)* [[FIELD_B]]
 //
 // Check accessors initialization
+// CHECK: [[KERNEL_BC:%[a-zA-Z0-9_]+]] = bitcast %union.__wrapper_union addrspace(4)* [[KERNEL_OBJ]] to %class.anon addrspace(4)*
+// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_BC]], i32 0, i32 0
+// CHECK: [[BITCAST:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP]] to %struct{{.*}}Base addrspace(4)*
 // CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST]], i32 0, i32 2
+// CHECK: [[ACC_FIELD_BC1:%[a-zA-Z0-9_]+]] = bitcast %"class.sycl::_V1::accessor" addrspace(4)* [[ACC_FIELD]] to i8 addrspace(4)*
+// CHECK: [[ACC_FIELD_BC2:%[a-zA-Z0-9_]+]] = bitcast i8 addrspace(4)* [[ACC_FIELD_BC1]] to %"class.sycl::_V1::accessor" addrspace(4)*
 // Default constructor call
-// CHECK: call spir_func void @_ZN4sycl3_V18accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC1Ev(%"class.sycl::_V1::accessor" addrspace(4)* {{[^,]*}} [[ACC_FIELD]])
+// CHECK: call spir_func void @_ZN4sycl3_V18accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC1Ev(%"class.sycl::_V1::accessor" addrspace(4)* {{[^,]*}} [[ACC_FIELD_BC2]])
+// Check __init method calls
+// CHECK: [[KERNEL_BC:%[a-zA-Z0-9_]+]] = bitcast %union.__wrapper_union addrspace(4)* [[KERNEL_OBJ]] to %class.anon addrspace(4)*
+// CHECK: [[GEP2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_BC]], i32 0, i32 0
+// CHECK: [[BITCAST3:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP2]] to %struct{{.*}}Base addrspace(4)*
+// CHECK: [[ACC1_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST3]], i32 0, i32 2
+// CHECK: [[ACC1_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(4)* [[ACC1_DATA]].addr.ascast
+// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.sycl::_V1::accessor" addrspace(4)* {{[^,]*}} [[ACC1_FIELD]], i8 addrspace(1)* noundef [[ACC1_DATA_LOAD]]
+
+// CHECK: [[KERNEL_BC:%[a-zA-Z0-9_]+]] = bitcast %union.__wrapper_union addrspace(4)* [[KERNEL_OBJ]] to %class.anon addrspace(4)*
+// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_BC]], i32 0, i32 0
 // CHECK: [[BITCAST1:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP]] to i8 addrspace(4)*
-// CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8 addrspace(4)* [[BITCAST1]], i64 20
+// CHECK: [[GEP1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds i8, i8 addrspace(4)* [[BITCAST1]], i64 20
 // CHECK: [[BITCAST2:%[a-zA-Z0-9_]+]] = bitcast i8 addrspace(4)* [[GEP1]] to %"class.sycl::_V1::accessor" addrspace(4)*
+// CHECK: [[ACC_FIELD_BC1:%[a-zA-Z0-9_]+]] = bitcast %"class.sycl::_V1::accessor" addrspace(4)* [[BITCAST2]] to i8 addrspace(4)*
+// CHECK: [[ACC_FIELD_BC2:%[a-zA-Z0-9_]+]] = bitcast i8 addrspace(4)* [[ACC_FIELD_BC1]] to %"class.sycl::_V1::accessor" addrspace(4)*
 // Default constructor call
-// CHECK: call spir_func void @_ZN4sycl3_V18accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC2Ev(%"class.sycl::_V1::accessor" addrspace(4)* {{[^,]*}} [[BITCAST2]])
+// CHECK: call spir_func void @_ZN4sycl3_V18accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC2Ev(%"class.sycl::_V1::accessor" addrspace(4)* {{[^,]*}} [[ACC_FIELD_BC2]])
 
 // CHECK C field initialization
 // CHECK: [[FIELD_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Captured, %struct{{.*}}Captured addrspace(4)* [[GEP]], i32 0, i32 2
diff --git a/clang/test/CodeGenSYCL/no_opaque_basic-kernel-wrapper.cpp b/clang/test/CodeGenSYCL/no_opaque_basic-kernel-wrapper.cpp
index ac5ee2ba871ab..55f949df1acce 100644
--- a/clang/test/CodeGenSYCL/no_opaque_basic-kernel-wrapper.cpp
+++ b/clang/test/CodeGenSYCL/no_opaque_basic-kernel-wrapper.cpp
@@ -27,12 +27,12 @@ int main() {
 // Check alloca for pointer argument
 // CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(1)*
 // Check lambda object alloca
-// CHECK: [[ANONALLOCA:%[a-zA-Z0-9_]+]] = alloca %class.anon
+// CHECK: [[UNIONALLOCA:%[a-zA-Z0-9_]+]] = alloca %union.__wrapper_union
 // Check allocas for ranges
 // CHECK: [[ARANGEA:%agg.tmp.*]] = alloca %"struct.sycl::_V1::range"
 // CHECK: [[MRANGEA:%agg.tmp.*]] = alloca %"struct.sycl::_V1::range"
 // CHECK: [[OIDA:%agg.tmp.*]] = alloca %"struct.sycl::_V1::id"
-// CHECK: [[ANON:%[a-zA-Z0-9_.]+]] = addrspacecast %class.anon* [[ANONALLOCA]] to %class.anon addrspace(4)*
+// CHECK: [[UNION:%[a-zA-Z0-9_.]+]] = addrspacecast %union.__wrapper_union* [[UNIONALLOCA]] to %union.__wrapper_union addrspace(4)*
 // CHECK: [[ARANGET:%agg.tmp.*]] = addrspacecast %"struct.sycl::_V1::range"* [[ARANGEA]] to %"struct.sycl::_V1::range" addrspace(4)*
 // CHECK: [[MRANGET:%agg.tmp.*]] = addrspacecast %"struct.sycl::_V1::range"* [[MRANGEA]] to %"struct.sycl::_V1::range" addrspace(4)*
 // CHECK: [[OIDT:%agg.tmp.*]] = addrspacecast %"struct.sycl::_V1::id"* [[OIDA]] to %"struct.sycl::_V1::id" addrspace(4)*
@@ -44,6 +44,7 @@ int main() {
 // CHECK: call spir_func {{.*}}accessor
 
 // Check accessor GEP
+// CHECK: [[ANON:%[a-zA-Z0-9_]+]] = bitcast %union.__wrapper_union addrspace(4)* [[UNION]] to %class.anon addrspace(4)*
 // CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* [[ANON]], i32 0, i32 0
 
 // Check load from kernel pointer argument alloca
diff --git a/clang/test/CodeGenSYCL/no_opaque_check-direct-attribute-propagation.cpp b/clang/test/CodeGenSYCL/no_opaque_check-direct-attribute-propagation.cpp
index 64777e1ba2202..d4ec6a37fbccc 100644
--- a/clang/test/CodeGenSYCL/no_opaque_check-direct-attribute-propagation.cpp
+++ b/clang/test/CodeGenSYCL/no_opaque_check-direct-attribute-propagation.cpp
@@ -317,7 +317,7 @@ int main() {
 
     // Test attribute is not propagated.
     // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name32() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
-    // CHECK: define {{.*}}spir_func void @{{.*}}Functor10{{.*}}(%class.Functor10 addrspace(4)* noundef align 1 dereferenceable_or_null(1) %this) #3 comdat align 2
+    // CHECK: define {{.*}}spir_func void @{{.*}}Functor10{{.*}}(%class.Functor10 addrspace(4)* noundef align 1 dereferenceable_or_null(1) %this) #2 comdat align 2
     // CHECK-NOT: noalias
     // CHECK-SAME: {
     // CHECK: define dso_local spir_func void @_Z4foo8v()
@@ -325,12 +325,12 @@ int main() {
     h.single_task<class kernel_name32>(f10);
 
     // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name33() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
-    // CHECK: define {{.*}}spir_func void @{{.*}}Foo8{{.*}}(%class.Foo8 addrspace(4)* noalias noundef align 1 dereferenceable_or_null(1) %this) #3 comdat align 2
+    // CHECK: define {{.*}}spir_func void @{{.*}}Foo8{{.*}}(%class.Foo8 addrspace(4)* noalias noundef align 1 dereferenceable_or_null(1) %this) #2 comdat align 2
     Foo8 boo8;
     h.single_task<class kernel_name33>(boo8);
 
     // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name34() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
-    // CHECK: define {{.*}}spir_func void @{{.*}}(%class.anon{{.*}} addrspace(4)* noalias noundef align 1 dereferenceable_or_null(1) %this) #4 align 2
+    // CHECK: define {{.*}}spir_func void @{{.*}}(%class.anon{{.*}} addrspace(4)* noalias noundef align 1 dereferenceable_or_null(1) %this) #3 align 2
     h.single_task<class kernel_name34>(
         []() [[intel::kernel_args_restrict]]{});
 
diff --git a/clang/test/CodeGenSYCL/no_opaque_image_accessor.cpp b/clang/test/CodeGenSYCL/no_opaque_image_accessor.cpp
index b181703395863..144b2df8db99e 100644
--- a/clang/test/CodeGenSYCL/no_opaque_image_accessor.cpp
+++ b/clang/test/CodeGenSYCL/no_opaque_image_accessor.cpp
@@ -8,27 +8,27 @@
 //
 // CHECK-1DRO: %opencl.image1d_ro_t = type opaque
 // CHECK-1DRO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image1d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
-// CHECK-1DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}sycl::_V1::accessor{{.*}} %{{[a-zA-Z]+}}, %opencl.image1d_ro_t addrspace(1)* %{{[0-9]+}})
-//
+// CHECK-1DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}sycl::_V1::accessor{{.*}} %{{[a-zA-Z0-9]+}}, %opencl.image1d_ro_t addrspace(1)* %{{[0-9]+}})
+
 // CHECK-2DRO: %opencl.image2d_ro_t = type opaque
 // CHECK-2DRO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image2d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
-// CHECK-2DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}sycl::_V1::accessor{{.*}} %{{[a-zA-Z]+}}, %opencl.image2d_ro_t addrspace(1)* %{{[0-9]+}})
+// CHECK-2DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}sycl::_V1::accessor{{.*}} %{{[a-zA-Z0-9]+}}, %opencl.image2d_ro_t addrspace(1)* %{{[0-9]+}})
 //
 // CHECK-3DRO: %opencl.image3d_ro_t = type opaque
 // CHECK-3DRO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image3d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
-// CHECK-3DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}sycl::_V1::accessor{{.*}} %{{[a-zA-Z]+}}, %opencl.image3d_ro_t addrspace(1)* %{{[0-9]+}})
+// CHECK-3DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}sycl::_V1::accessor{{.*}} %{{[a-zA-Z0-9]+}}, %opencl.image3d_ro_t addrspace(1)* %{{[0-9]+}})
 //
 // CHECK-1DWO: %opencl.image1d_wo_t = type opaque
 // CHECK-1DWO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image1d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
-// CHECK-1DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}sycl::_V1::accessor{{.*}} %{{[a-zA-Z]+}}, %opencl.image1d_wo_t addrspace(1)* %{{[0-9]+}})
+// CHECK-1DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}sycl::_V1::accessor{{.*}} %{{[a-zA-Z0-9]+}}, %opencl.image1d_wo_t addrspace(1)* %{{[0-9]+}})
 //
 // CHECK-2DWO: %opencl.image2d_wo_t = type opaque
 // CHECK-2DWO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image2d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
-// CHECK-2DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}sycl::_V1::accessor{{.*}} %{{[a-zA-Z]+}}, %opencl.image2d_wo_t addrspace(1)* %{{[0-9]+}})
+// CHECK-2DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}sycl::_V1::accessor{{.*}} %{{[a-zA-Z0-9]+}}, %opencl.image2d_wo_t addrspace(1)* %{{[0-9]+}})
 //
 // CHECK-3DWO: %opencl.image3d_wo_t = type opaque
 // CHECK-3DWO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image3d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
-// CHECK-3DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}sycl::_V1::accessor{{.*}} %{{[a-zA-Z]+}}, %opencl.image3d_wo_t addrspace(1)* %{{[0-9]+}})
+// CHECK-3DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}sycl::_V1::accessor{{.*}} %{{[a-zA-Z0-9]+}}, %opencl.image3d_wo_t addrspace(1)* %{{[0-9]+}})
 //
 // TODO: Add tests for the image_array opencl datatype support.
 #include "Inputs/sycl.hpp"
diff --git a/clang/test/CodeGenSYCL/no_opaque_inheritance.cpp b/clang/test/CodeGenSYCL/no_opaque_inheritance.cpp
index 1f8bcb83d3de3..7c174ca2de08b 100644
--- a/clang/test/CodeGenSYCL/no_opaque_inheritance.cpp
+++ b/clang/test/CodeGenSYCL/no_opaque_inheritance.cpp
@@ -54,31 +54,33 @@ int main() {
 
 // Check allocas for kernel parameters and local functor object
 // CHECK: %[[ARG_A_ALLOCA:[a-zA-Z0-9_.]+]] = alloca i32, align 4
-// CHECK: %[[LOCAL_OBJECT_ALLOCA:[a-zA-Z0-9_.]+]] = alloca %struct.derived, align 8
+// CHECK: %[[LOCAL_OBJECT_ALLOCA:[a-zA-Z0-9_.]+]] = alloca %union.__wrapper_union, align 8
 // CHECK: %[[ARG_A:[a-zA-Z0-9_.]+]] = addrspacecast i32* %[[ARG_A_ALLOCA]] to i32 addrspace(4)*
-// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = addrspacecast %struct.derived* %[[LOCAL_OBJECT_ALLOCA]] to %struct.derived addrspace(4)*
+// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = addrspacecast %union.__wrapper_union* %[[LOCAL_OBJECT_ALLOCA]] to %union.__wrapper_union addrspace(4)*
 // CHECK: %[[ARG_BASE:[a-zA-Z0-9_.]+]] = addrspacecast %struct.base* %_arg__base to %struct.base addrspace(4)*
 // CHECK: %[[ARG_BASE1:[a-zA-Z0-9_.]+]] = addrspacecast %class.__generated_second_base* %_arg__base1 to %class.__generated_second_base addrspace(4)*
 // CHECK: store i32 %_arg_a, i32 addrspace(4)* %[[ARG_A]], align 4
 
 // Initialize 'base' subobject
-// CHECK: %[[DERIVED_TO_BASE:.*]] = bitcast %struct.derived addrspace(4)* %[[LOCAL_OBJECT]] to %struct.base addrspace(4)*
+// CHECK: %[[DERIVED_PTR:.*]] = bitcast %union.__wrapper_union addrspace(4)* %[[LOCAL_OBJECT]] to %struct.derived addrspace(4)*
+// CHECK: %[[DERIVED_TO_BASE:.*]] = bitcast %struct.derived addrspace(4)* %[[DERIVED_PTR]] to %struct.base addrspace(4)*
 // CHECK: %[[BASE_TO_PTR:.*]] = bitcast %struct.base addrspace(4)* %[[DERIVED_TO_BASE]] to i8 addrspace(4)*
 // CHECK: %[[PARAM_TO_PTR:.*]] = bitcast %struct.base addrspace(4)* %[[ARG_BASE]] to i8 addrspace(4)*
 // CHECK: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 8 %[[BASE_TO_PTR]], i8 addrspace(4)* 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 addrspace(4)* %[[LOCAL_OBJECT]] to i8 addrspace(4)*
-// CHECK: %[[OFFSET_CALC:.*]] = getelementptr inbounds i8, i8 addrspace(4)* %[[DERIVED_PTR]], i64 16
+// CHECK: %[[DERIVED_PTR:.*]] = bitcast %union.__wrapper_union addrspace(4)* %[[LOCAL_OBJECT]] to %struct.derived addrspace(4)*
+// CHECK: %[[BASE_PTR:.*]] = bitcast %struct.derived addrspace(4)* %[[DERIVED_PTR]] to i8 addrspace(4)*
+// CHECK: %[[OFFSET_CALC:.*]] = getelementptr inbounds i8, i8 addrspace(4)* %[[BASE_PTR]], i64 16
 // CHECK: %[[TO_SECOND_BASE:.*]] = bitcast i8 addrspace(4)* %[[OFFSET_CALC]] to %class.second_base addrspace(4)*
-// CHECK: %[[GEN_TO_SECOND_BASE:.*]] = bitcast %class.__generated_second_base addrspace(4)* %[[ARG_BASE1]] to %class.second_base addrspace(4)*
 // CHECK: %[[TO:.*]] = bitcast %class.second_base addrspace(4)* %[[TO_SECOND_BASE]] to i8 addrspace(4)*
-// CHECK: %[[FROM:.*]] = bitcast %class.second_base addrspace(4)* %[[GEN_TO_SECOND_BASE]] to i8 addrspace(4)*
+// CHECK: %[[FROM:.*]] = bitcast %class.__generated_second_base addrspace(4)* %[[ARG_BASE1]] to i8 addrspace(4)*
 // CHECK: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 8 %[[TO]], i8 addrspace(4)* align 8 %[[FROM]], i64 8, i1 false)
 
 
 // Initialize field 'a'
-// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.derived, %struct.derived addrspace(4)* %[[LOCAL_OBJECT]], i32 0, i32 3
 // CHECK: %[[LOAD_A:[0-9]+]] = load i32, i32 addrspace(4)* %[[ARG_A]], align 4
+// CHECK: %[[DERIVED_PTR:.*]] = bitcast %union.__wrapper_union addrspace(4)* %[[LOCAL_OBJECT]] to %struct.derived addrspace(4)*
+// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.derived, %struct.derived addrspace(4)* %[[DERIVED_PTR]], i32 0, i32 3
 // CHECK: store i32 %[[LOAD_A]], i32 addrspace(4)* %[[GEP_A]]
diff --git a/clang/test/CodeGenSYCL/no_opaque_kernel-param-acc-array.cpp b/clang/test/CodeGenSYCL/no_opaque_kernel-param-acc-array.cpp
index 10ed0a1f268f7..e51d4e3423b45 100644
--- a/clang/test/CodeGenSYCL/no_opaque_kernel-param-acc-array.cpp
+++ b/clang/test/CodeGenSYCL/no_opaque_kernel-param-acc-array.cpp
@@ -39,7 +39,7 @@ int main() {
 // CHECK: [[MEM_ARG2:%[a-zA-Z0-9_.]+]] = alloca i32 addrspace(1)*, align 8
 
 // CHECK lambda object alloca
-// CHECK: [[LOCAL_OBJECTA:%__SYCLKernel]] = alloca %class.anon, align 4
+// CHECK: [[LOCAL_OBJECTA:%__wrapper_union]] = alloca %union.__wrapper_union, align 4
 
 // CHECK allocas for ranges
 // CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::range"
diff --git a/clang/test/CodeGenSYCL/no_opaque_sampler.cpp b/clang/test/CodeGenSYCL/no_opaque_sampler.cpp
index ba0d9d37a8487..74055e613346a 100644
--- a/clang/test/CodeGenSYCL/no_opaque_sampler.cpp
+++ b/clang/test/CodeGenSYCL/no_opaque_sampler.cpp
@@ -2,14 +2,45 @@
 // CHECK: define {{.*}}spir_kernel void @{{[a-zA-Z0-9_]+}}(%opencl.sampler_t addrspace(2)* [[SAMPLER_ARG:%[a-zA-Z0-9_]+]])
 // CHECK-NEXT: entry:
 // CHECK-NEXT: [[SAMPLER_ARG]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8
-// CHECK: [[ANON:%[a-zA-Z0-9_]+]] = alloca %class.anon, align 8
-// CHECK: [[ANONCAST:%[a-zA-Z0-9_.]+]] = addrspacecast %class.anon* [[ANON]] to %class.anon addrspace(4)*
+// CHECK: [[ANON:%[a-zA-Z0-9_]+]] = alloca %union.__wrapper_union, align 8
+// CHECK: [[ANONCAST:%[a-zA-Z0-9_.]+]] = addrspacecast %union.__wrapper_union* [[ANON]] to %union.__wrapper_union addrspace(4)*
 // CHECK: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG]], %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG]].addr.ascast, align 8
-// CHECK-NEXT: [[GEP:%[a-zA-z0-9]+]]  = getelementptr inbounds %class.anon, %class.anon addrspace(4)* [[ANONCAST]], i32 0, i32 0
+// CHECK-NEXT: [[BITCAST:%[0-9]+]] = bitcast %union.__wrapper_union* [[ANON]] to i8*
+// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[BITCAST]])
+// CHECK: [[GEP:%[a-zA-z0-9]+]]  = getelementptr inbounds %union.__wrapper_union, %union.__wrapper_union addrspace(4)* [[ANONCAST]], i32 0, i32 0
 // CHECK-NEXT: [[LOAD_SAMPLER_ARG:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG]].addr.ascast, align 8
 // CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.sycl::_V1::sampler" addrspace(4)* {{[^,]*}} [[GEP]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]])
 //
 
+//   %_arg_smplr.addr = alloca %opencl.sampler_t addrspace(2)*, align 8
+//   %__wrapper_union = alloca %union.__wrapper_union, align 8
+//   %0 = alloca %class.anon addrspace(4)*, align 8
+//   %_arg_smplr.addr.ascast = addrspacecast %opencl.sampler_t addrspace(2)** %_arg_smplr.addr to %opencl.sampler_t addrspace(2)* addrspace(4)*
+//   %__wrapper_union.ascast = addrspacecast %union.__wrapper_union* %__wrapper_union to %union.__wrapper_union addrspace(4)*
+//   %1 = addrspacecast %class.anon addrspace(4)** %0 to %class.anon addrspace(4)* addrspace(4)*
+//   store %opencl.sampler_t addrspace(2)* %_arg_smplr, %opencl.sampler_t addrspace(2)* addrspace(4)* %_arg_smplr.addr.ascast, align 8, !tbaa !14
+//   %2 = bitcast %union.__wrapper_union* %__wrapper_union to i8*
+//   call void @llvm.lifetime.start.p0i8(i64 8, i8* %2) #5
+//   %3 = bitcast %class.anon addrspace(4)** %0 to i8*
+//   call void @llvm.lifetime.start.p0i8(i64 8, i8* %3) #5
+//   %4 = bitcast %union.__wrapper_union addrspace(4)* %__wrapper_union.ascast to %class.anon addrspace(4)*
+//   store %class.anon addrspace(4)* %4, %class.anon addrspace(4)* addrspace(4)* %1, align 8, !tbaa !18
+//   %5 = load %class.anon addrspace(4)*, %class.anon addrspace(4)* addrspace(4)* %1, align 8, !tbaa !18
+//   %smplr = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %5, i32 0, i32 0
+//   %6 = bitcast %"class.sycl::_V1::sampler" addrspace(4)* %smplr to i8 addrspace(4)*
+//   %7 = bitcast i8 addrspace(4)* %6 to %"class.sycl::_V1::sampler" addrspace(4)*
+//   %8 = load %class.anon addrspace(4)*, %class.anon addrspace(4)* addrspace(4)* %1, align 8, !tbaa !18
+//   %smplr1 = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %8, i32 0, i32 0
+//   %9 = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)* addrspace(4)* %_arg_smplr.addr.ascast, align 8, !tbaa !14
+//   call spir_func void @_ZN4sycl3_V17sampler6__initE11ocl_sampler(%"class.sycl::_V1::sampler" addrspace(4)* noundef align 8 dereferenceable_or_null(8) %smplr1, %opencl.sampler_t addrspace(2)* %9) #6
+//   %10 = load %class.anon addrspace(4)*, %class.anon addrspace(4)* addrspace(4)* %1, align 8, !tbaa !18
+//   call spir_func void @_ZZ4mainENKUlvE_clEv(%class.anon addrspace(4)* noundef align 8 dereferenceable_or_null(8) %10) #6
+//   %11 = bitcast %class.anon addrspace(4)** %0 to i8*
+//   call void @llvm.lifetime.end.p0i8(i64 8, i8* %11) #5
+//   %12 = bitcast %union.__wrapper_union* %__wrapper_union to i8*
+//   call void @llvm.lifetime.end.p0i8(i64 8, i8* %12) #5
+//   ret void
+
 // CHECK: define {{.*}}spir_kernel void @{{[a-zA-Z0-9_]+}}(%opencl.sampler_t addrspace(2)* [[SAMPLER_ARG_WRAPPED:%[a-zA-Z0-9_]+]], i32 noundef [[ARG_A:%[a-zA-Z0-9_]+]])
 
 // Check alloca
diff --git a/clang/test/CodeGenSYCL/no_opaque_stall_enable_device.cpp b/clang/test/CodeGenSYCL/no_opaque_stall_enable_device.cpp
index 8d3a426977288..99751c4be681f 100644
--- a/clang/test/CodeGenSYCL/no_opaque_stall_enable_device.cpp
+++ b/clang/test/CodeGenSYCL/no_opaque_stall_enable_device.cpp
@@ -26,12 +26,12 @@ class Foo {
 int main() {
   q.submit([&](handler &h) {
     // CHECK: define {{.*}}spir_kernel void @{{.*}}test_kernel1() {{.*}} !stall_enable ![[NUM4:[0-9]+]]
-    // CHECK: define {{.*}}spir_func void @{{.*}}FuncObjclEv(%struct.{{.*}}FuncObj addrspace(4)* noundef align 1 dereferenceable_or_null(1) %this) #2 comdat align 2{{.*}} !stall_enable ![[NUM4]]
+    // CHECK: define {{.*}}spir_func void @{{.*}}FuncObjclEv(%struct.{{.*}}FuncObj addrspace(4)* noundef align 1 dereferenceable_or_null(1) %this) #1 comdat align 2{{.*}} !stall_enable ![[NUM4]]
     h.single_task<class test_kernel1>(
         FuncObj());
 
     // CHECK: define {{.*}}spir_kernel void @{{.*}}test_kernel2() {{.*}} !stall_enable ![[NUM4]]
-    // CHECK define {{.*}}spir_func void @{{.*}}FooclEv(%class._ZTS3Foo.Foo addrspace(4)* noundef align 1 dereferenceable_or_null(1) %this) #2 comdat align 2{{.*}} !stall_enable ![[NUM4]]
+    // CHECK define {{.*}}spir_func void @{{.*}}FooclEv(%class._ZTS3Foo.Foo addrspace(4)* noundef align 1 dereferenceable_or_null(1) %this) #1 comdat align 2{{.*}} !stall_enable ![[NUM4]]
     Foo f;
     h.single_task<class test_kernel2>(f);
 
@@ -47,7 +47,7 @@ int main() {
     // CHECK: define {{.*}}spir_kernel void @{{.*}}test_kernel4()
     // CHECK-NOT: !stall_enable
     // CHECK-SAME: {
-    // CHECK: define {{.*}}spir_func void @{{.*}}func1{{.*}}(%class.anon{{.*}} addrspace(4)* noundef align 1 dereferenceable_or_null(1) %this) #2 align 2{{.*}} !stall_enable ![[NUM4]]
+    // CHECK: define {{.*}}spir_func void @{{.*}}func1{{.*}}(%class.anon{{.*}} addrspace(4)* noundef align 1 dereferenceable_or_null(1) %this) #1 align 2{{.*}} !stall_enable ![[NUM4]]
     h.single_task<class test_kernel4>(
         []() { func1(); });
 
diff --git a/clang/test/CodeGenSYCL/no_opaque_union-kernel-param.cpp b/clang/test/CodeGenSYCL/no_opaque_union-kernel-param.cpp
index a9092aae39e91..735a2e0ec0b47 100644
--- a/clang/test/CodeGenSYCL/no_opaque_union-kernel-param.cpp
+++ b/clang/test/CodeGenSYCL/no_opaque_union-kernel-param.cpp
@@ -31,12 +31,16 @@ int main() {
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A(%union.MyUnion* noundef byval(%union.MyUnion) align 4 [[MEM_ARG:%[a-zA-Z0-9_]+]])
 
 // Check lambda object alloca
-// CHECK: [[LOCAL_OBJECT:%__SYCLKernel]] = alloca %class.anon, align 4
+// CHECK: [[LOCAL_OBJECT:%__wrapper_union]] = alloca %union.__wrapper_union, align 4
+// CHECK: [[FUNCTOR_PTR:%.*]] = alloca %class.anon addrspace(4)*
 
-// CHECK: [[LOCAL_OBJECTAS:%.*]] = addrspacecast %class.anon* [[LOCAL_OBJECT]] to %class.anon addrspace(4)*
+// CHECK: [[LOCAL_OBJECTAS:%.*]] = addrspacecast %union.__wrapper_union* [[LOCAL_OBJECT]] to %union.__wrapper_union addrspace(4)*
+// CHECK: [[FUNCTOR_PTRAS:%.*]] = addrspacecast %class.anon addrspace(4)** [[FUNCTOR_PTR]] to %class.anon addrspace(4)* addrspace(4)*
 // CHECK: [[MEM_ARGAS:%.*]] = addrspacecast %union.MyUnion* [[MEM_ARG]] to %union.MyUnion addrspace(4)*
-// CHECK: [[L_STRUCT_ADDR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* [[LOCAL_OBJECTAS]], i32 0, i32 0
+// CHECK: [[ANON_OBJECTAS:%.*]] = bitcast %union.__wrapper_union addrspace(4)* [[LOCAL_OBJECTAS]] to %class.anon addrspace(4)*
+// CHECK: [[L_STRUCT_ADDR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* [[ANON_OBJECTAS]], i32 0, i32 0
 // CHECK: [[MEMCPY_DST:%[0-9a-zA-Z_]+]] = bitcast %union.{{.*}}MyUnion addrspace(4)* [[L_STRUCT_ADDR]] to i8 addrspace(4)*
 // CHECK: [[MEMCPY_SRC:%[0-9a-zA-Z_]+]] = bitcast %union.{{.*}}MyUnion addrspace(4)* [[MEM_ARGAS]] to i8 addrspace(4)*
 // CHECK: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 4 [[MEMCPY_DST]], i8 addrspace(4)* align 4 [[MEMCPY_SRC]], i64 12, i1 false)
-// CHECK: call spir_func void @{{.*}}(%class.anon addrspace(4)* {{[^,]*}} [[LOCAL_OBJECTAS]])
+// CHECK: [[FUNCTOR:%[0-9a-zA-Z_]+]] = load %class.anon addrspace(4)*, %class.anon addrspace(4)* addrspace(4)* [[FUNCTOR_PTRAS]]
+// CHECK: call spir_func void @{{.*}}(%class.anon addrspace(4)* {{[^,]*}} [[FUNCTOR]])
diff --git a/clang/test/CodeGenSYCL/stall_enable_device.cpp b/clang/test/CodeGenSYCL/stall_enable_device.cpp
index 1cdfd4ee4c189..86de0d8c53b3f 100644
--- a/clang/test/CodeGenSYCL/stall_enable_device.cpp
+++ b/clang/test/CodeGenSYCL/stall_enable_device.cpp
@@ -26,12 +26,12 @@ class Foo {
 int main() {
   q.submit([&](handler &h) {
     // CHECK: define {{.*}}spir_kernel void @{{.*}}test_kernel1() {{.*}} !stall_enable ![[NUM4:[0-9]+]]
-    // CHECK: define {{.*}}spir_func void @{{.*}}FuncObjclEv(ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) %this) #2 comdat align 2{{.*}} !stall_enable ![[NUM4]]
+    // CHECK: define {{.*}}spir_func void @{{.*}}FuncObjclEv(ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) %this) #1 comdat align 2{{.*}} !stall_enable ![[NUM4]]
     h.single_task<class test_kernel1>(
         FuncObj());
 
     // CHECK: define {{.*}}spir_kernel void @{{.*}}test_kernel2() {{.*}} !stall_enable ![[NUM4]]
-    // CHECK define {{.*}}spir_func void @{{.*}}FooclEv(ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) %this) #2 comdat align 2{{.*}} !stall_enable ![[NUM4]]
+    // CHECK define {{.*}}spir_func void @{{.*}}FooclEv(ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) %this) #1 comdat align 2{{.*}} !stall_enable ![[NUM4]]
     Foo f;
     h.single_task<class test_kernel2>(f);
 
@@ -47,7 +47,7 @@ int main() {
     // CHECK: define {{.*}}spir_kernel void @{{.*}}test_kernel4()
     // CHECK-NOT: !stall_enable
     // CHECK-SAME: {
-    // CHECK: define {{.*}}spir_func void @{{.*}}func1{{.*}}(ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) %this) #2 align 2{{.*}} !stall_enable ![[NUM4]]
+    // CHECK: define {{.*}}spir_func void @{{.*}}func1{{.*}}(ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) %this) #1 align 2{{.*}} !stall_enable ![[NUM4]]
     h.single_task<class test_kernel4>(
         []() { func1(); });
 
diff --git a/clang/test/CodeGenSYCL/union-kernel-param.cpp b/clang/test/CodeGenSYCL/union-kernel-param.cpp
index 0cd22146f5f9d..2f485ffa5b205 100644
--- a/clang/test/CodeGenSYCL/union-kernel-param.cpp
+++ b/clang/test/CodeGenSYCL/union-kernel-param.cpp
@@ -27,10 +27,12 @@ int main() {
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A(ptr noundef byval(%union.MyUnion) align 4 [[MEM_ARG:%[a-zA-Z0-9_]+]])
 
 // Check lambda object alloca
-// CHECK: [[LOCAL_OBJECT:%__SYCLKernel]] = alloca %class.anon, align 4
+// CHECK: [[LOCAL_OBJECT:%__wrapper_union]] = alloca %union.__wrapper_union, align 4
 
 // CHECK: [[LOCAL_OBJECTAS:%.*]] = addrspacecast ptr [[LOCAL_OBJECT]] to ptr addrspace(4)
 // CHECK: [[MEM_ARGAS:%.*]] = addrspacecast ptr [[MEM_ARG]] to ptr addrspace(4)
 // CHECK: [[L_STRUCT_ADDR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, ptr addrspace(4) [[LOCAL_OBJECTAS]], i32 0, i32 0
 // CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[L_STRUCT_ADDR]], ptr addrspace(4) align 4 [[MEM_ARGAS]], i64 12, i1 false)
-// CHECK: call spir_func void @{{.*}}(ptr addrspace(4) {{[^,]*}} [[LOCAL_OBJECTAS]])
+// CHECK: store ptr addrspace(4) [[LOCAL_OBJECTAS]], ptr addrspace(4) [[KERNEL_REF_ADDR:%[A-Za-z0-9]*]]
+// CHECK: [[KERNEL_REF:%[A-Za-z0-9]*]] = load ptr addrspace(4), ptr addrspace(4) [[KERNEL_REF_ADDR]]
+// CHECK: call spir_func void @{{.*}}(ptr addrspace(4) {{[^,]*}} [[KERNEL_REF]])
diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp
index a1ef7bf504641..c110f28872804 100644
--- a/clang/test/SemaSYCL/Inputs/sycl.hpp
+++ b/clang/test/SemaSYCL/Inputs/sycl.hpp
@@ -1,6 +1,8 @@
 #ifndef SYCL_HPP
 #define SYCL_HPP
 
+void* operator new  (__SIZE_TYPE__ size, void* ptr) noexcept;
+void* operator new[](__SIZE_TYPE__ size, void* ptr) noexcept;
 #define __SYCL_TYPE(x) [[__sycl_detail__::sycl_type(x)]]
 
 // Shared code for SYCL tests
diff --git a/clang/test/SemaSYCL/accessor_inheritance.cpp b/clang/test/SemaSYCL/accessor_inheritance.cpp
index a631116466896..345a4e9c600eb 100644
--- a/clang/test/SemaSYCL/accessor_inheritance.cpp
+++ b/clang/test/SemaSYCL/accessor_inheritance.cpp
@@ -44,32 +44,97 @@ int main() {
 // CHECK: ParmVarDecl{{.*}} used _arg_C 'int'
 
 // Check lambda initialization
-// CHECK: VarDecl {{.*}} used __SYCLKernel '(lambda at {{.*}}accessor_inheritance.cpp
-// CHECK-NEXT: InitListExpr {{.*}}
-// CHECK-NEXT: InitListExpr {{.*}} 'AccessorDerived'
-// CHECK-NEXT: InitListExpr {{.*}} 'AccessorBase'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_A' 'int'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_B' 'int'
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::accessor<char, 1, sycl::access::mode::read>':'sycl::accessor<char, 1, sycl::access::mode::read, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>' 'void () noexcept'
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::accessor<char, 1, sycl::access::mode::read>':'sycl::accessor<char, 1, sycl::access::mode::read, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>' 'void () noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_C' 'int'
+// CHECK: VarDecl {{.*}} used __wrapper_union '__wrapper_union'
+// CHECK: VarDecl {{.*}} used '(lambda at {{.*}}accessor_inheritance.cpp{{.*}} &'
+// CHECK-NEXT: MemberExpr {{.*}} '(lambda at {{.*}}accessor_inheritance.cpp{{.*}})' lvalue . {{.*}}
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_union' lvalue Var {{.*}} '__wrapper_union' '__wrapper_union'
+// Init A
+// memcpy(lambda.A, __arg_A, 4)
+//
+// CHECK: CallExpr {{.*}} 'void *'
+// CHECK: DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK: UnaryOperator {{.*}} '&'
+// CHECK:  MemberExpr {{.*}} 'int' lvalue .A
+// CHECK:   ImplicitCastExpr {{.*}} <DerivedToBase (AccessorBase)>
+// CHECK:    MemberExpr {{.*}} 'AccessorDerived':'AccessorDerived' lvalue .
+// CHECK:     DeclRefExpr {{.*}} Var {{.*}} '' '(lambda at {{.*}}accessor_inheritance.cpp{{.*}}) &'
+// CHECK: UnaryOperator {{.*}} prefix '&'
+// CHECK:  DeclRefExpr {{.*}} '_arg_A'
+// CHECK: IntegerLiteral {{.*}} 'unsigned long' 4
 
-// Check __init calls
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} .__init
-// CHECK-NEXT: MemberExpr {{.*}} .AccField
-// CHECK-NEXT: ImplicitCastExpr {{.*}}'AccessorBase' lvalue <DerivedToBase (AccessorBase)>
-// CHECK-NEXT: MemberExpr {{.*}}'AccessorDerived' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}}'(lambda at {{.*}}accessor_inheritance.cpp
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global char *' <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} '__global char *' lvalue ParmVar {{.*}} '_arg_AccField' '__global char *'
+// Init B
+// memcpy(lambda.B, __arg_B, 4)
+//
+// CHECK: CallExpr {{.*}} 'void *'
+// CHECK: DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK: UnaryOperator {{.*}} '&'
+// CHECK:  MemberExpr {{.*}} 'int' lvalue .B
+// CHECK:   ImplicitCastExpr {{.*}} <DerivedToBase (AccessorBase)>
+// CHECK:    MemberExpr {{.*}} 'AccessorDerived':'AccessorDerived' lvalue .
+// CHECK:     DeclRefExpr {{.*}} Var {{.*}} '' '(lambda at {{.*}}accessor_inheritance.cpp{{.*}}) &'
+// CHECK: UnaryOperator {{.*}} prefix '&'
+// CHECK:  DeclRefExpr {{.*}} '_arg_B'
+// CHECK: IntegerLiteral {{.*}} 'unsigned long' 4
 
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr{{.*}} lvalue .__init
-// CHECK-NEXT: MemberExpr{{.*}}'AccessorDerived' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}accessor_inheritance.cpp
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global char *' <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} '__global char *' lvalue ParmVar {{.*}} '_arg__base' '__global char *'
+// Init AccField
+// placement new
+//
+// CHECK: CXXNewExpr {{.*}} 'sycl::accessor<char, 1, sycl::access::mode::read, sycl::access::target::global_buffer, sycl::access::placeholder::false_t> *' global Function {{.*}} 'operator new'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::accessor<char, 1, sycl::access::mode::read, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>' 'void () noexcept'
+// CHECK: MemberExpr {{.*}}'sycl::accessor<char, 1, sycl::access::mode::read, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>' lvalue .AccField
+// CHECK-NEXT: ImplicitCastExpr {{.*}} <DerivedToBase (AccessorBase)>
+// CHECK-NEXT: MemberExpr {{.*}} 'AccessorDerived':'AccessorDerived'
+// CHECK-NEXT: DeclRefExpr {{.*}} Var {{.*}} '' '(lambda at {{.*}}accessor_inheritance.cpp{{.*}}) &'
+
+// call to __init
+//
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
+// CHECK-NEXT: MemberExpr {{.*}} 'void (PtrType, range<1>, range<1>, id<1>)' lvalue .__init
+// CHECK-NEXT:  MemberExpr {{.*}} .AccField {{.*}}
+// CHECK-NEXT:   ImplicitCastExpr {{.*}} <DerivedToBase (AccessorBase)>
+// CHECK-NEXT:   MemberExpr {{.*}} 'AccessorDerived':'AccessorDerived'
+// CHECK-NEXT:   DeclRefExpr {{.*}} Var {{.*}} '' '(lambda at {{.*}}accessor_inheritance.cpp{{.*}}) &'
+// CHECK:       DeclRefExpr {{.*}} '_arg_AccField' '__global char *'
+// CHECK-NEXT:  CXXConstructExpr {{.*}} 'range<1>':'sycl::range<1>'
+// CHECK:        DeclRefExpr {{.*}} ParmVar {{.*}} '_arg_AccField' 'sycl::range<1>'
+// CHECK-NEXT:  CXXConstructExpr {{.*}} 'range<1>':'sycl::range<1>'
+// CHECK:        DeclRefExpr {{.*}} ParmVar {{.*}} '_arg_AccField' 'sycl::range<1>'
+// CHECK-NEXT:  CXXConstructExpr {{.*}} 'id<1>':'sycl::id<1>'
+// CHECK:        DeclRefExpr {{.*}} ParmVar {{.*}} '_arg_AccField' 'sycl::id<1>'
+
+// Init inherited accessor
+// placement new
+//
+// CHECK:      CXXNewExpr {{.*}} 'sycl::accessor<char, 1, sycl::access::mode::read, sycl::access::target::global_buffer, sycl::access::placeholder::false_t> *' global Function {{.*}} 'operator new'
+// CHECK-NEXT:  CXXConstructExpr {{.*}} 'sycl::accessor<char, 1, sycl::access::mode::read, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>' 'void () noexcept'
+// CHECK:       ImplicitCastExpr {{.*}} <DerivedToBase (accessor)>
+// CHECK-NEXT:   MemberExpr {{.*}} 'AccessorDerived':'AccessorDerived'
+// CHECK-NEXT:    DeclRefExpr {{.*}} Var {{.*}} '' '(lambda at {{.*}}accessor_inheritance.cpp{{.*}}) &'
+
+// call to __init
+//
+// CHECK:      CXXMemberCallExpr {{.*}} 'void'
+// CHECK-NEXT: MemberExpr {{.*}} 'void (PtrType, range<1>, range<1>, id<1>)' lvalue .__init
+// CHECK-NEXT:  ImplicitCastExpr {{.*}} <DerivedToBase (accessor)>
+// CHECK-NEXT:   MemberExpr {{.*}} 'AccessorDerived':'AccessorDerived'
+// CHECK-NEXT:   DeclRefExpr {{.*}} Var {{.*}} '' '(lambda at {{.*}}accessor_inheritance.cpp{{.*}}) &'
+// CHECK:       DeclRefExpr {{.*}} '_arg__base' '__global char *'
+// CHECK-NEXT:  CXXConstructExpr {{.*}} 'range<1>':'sycl::range<1>'
+// CHECK:        DeclRefExpr {{.*}} ParmVar {{.*}} '_arg__base' 'sycl::range<1>'
+// CHECK-NEXT:  CXXConstructExpr {{.*}} 'range<1>':'sycl::range<1>'
+// CHECK:        DeclRefExpr {{.*}} ParmVar {{.*}} '_arg__base' 'sycl::range<1>'
+// CHECK-NEXT:  CXXConstructExpr {{.*}} 'id<1>':'sycl::id<1>'
+// CHECK:        DeclRefExpr {{.*}} ParmVar {{.*}} '_arg__base' 'sycl::id<1>'
+
+// Init C
+// memcpy(lambda.C, __arg_C, 4)
+//
+// CHECK: CallExpr {{.*}} 'void *'
+// CHECK: DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK:  UnaryOperator {{.*}} '&'
+// CHECK:  MemberExpr {{.*}} 'int' lvalue .C
+// CHECK:  MemberExpr {{.*}} 'AccessorDerived':'AccessorDerived' lvalue .
+// CHECK:  DeclRefExpr {{.*}} Var {{.*}} '' '(lambda at {{.*}}accessor_inheritance.cpp{{.*}}) &'
+// CHECK: UnaryOperator {{.*}} prefix '&'
+// CHECK:  DeclRefExpr {{.*}} '_arg_C'
+// CHECK: IntegerLiteral {{.*}} 'unsigned long' 4
diff --git a/clang/test/SemaSYCL/array-kernel-param.cpp b/clang/test/SemaSYCL/array-kernel-param.cpp
index 221aa855fa1f6..edec40337fbd9 100644
--- a/clang/test/SemaSYCL/array-kernel-param.cpp
+++ b/clang/test/SemaSYCL/array-kernel-param.cpp
@@ -136,18 +136,24 @@ int main() {
 // Check Kernel_Array inits
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
-// CHECK-NEXT: VarDecl {{.*}} cinit
-// CHECK-NEXT: InitListExpr
-// 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_Array' '__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_Array' '__wrapper_class'
+// CHECK-NEXT:  VarDecl
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT:  VarDecl {{.*}} cinit
+// CHECK-NEXT:   MemberExpr
+// CHECK-NEXT:    DeclRefExpr
+// CHECK-NEXT: CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .Array
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT:  UnaryOperator
+// CHECK-NEXT:   MemberExpr
+// CHECK-NEXT:    DeclRefExpr {{.*}} ParmVar {{.*}} '_arg_Array'
+// CHECK-NEXT: IntegerLiteral {{.*}} 8
+
 
 // Check Kernel_Array_Ptrs parameters
 // CHECK: FunctionDecl {{.*}}Kernel_Array_Ptrs{{.*}} 'void (__wrapper_class)'
@@ -155,29 +161,23 @@ int main() {
 // Check Kernel_Array_Ptrs inits
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
-// CHECK-NEXT: VarDecl {{.*}} cinit
-// CHECK-NEXT: InitListExpr
-// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]'
-// Initializer for ArrayOfPointers[0]
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
-// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
-// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
-// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers'
-// CHECK-NEXT: IntegerLiteral {{.*}} 0
-// Initializer for ArrayOfPointers[1]
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
-// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
-// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
-// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers'
-// CHECK-NEXT: IntegerLiteral {{.*}} 1
+// CHECK-NEXT:  VarDecl
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT:  VarDecl {{.*}} cinit
+// CHECK-NEXT:   MemberExpr
+// CHECK-NEXT:    DeclRefExpr
+// CHECK-NEXT: CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .Array
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT:  UnaryOperator
+// CHECK-NEXT:   MemberExpr
+// CHECK-NEXT:    DeclRefExpr {{.*}} ParmVar {{.*}} '_arg_ArrayOfPointers'
+// CHECK-NEXT: IntegerLiteral {{.*}} 16
 
 // Check Kernel_StructAccArray parameters
 // CHECK: FunctionDecl {{.*}}Kernel_StructAccArray{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
@@ -193,29 +193,59 @@ int main() {
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'sycl::id<1>'
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
-// CHECK-NEXT: VarDecl {{.*}} used __SYCLKernel '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' cinit
-// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp{{.*}})'
-// CHECK-NEXT: InitListExpr {{.*}} 'StructWithAccessors'
-// CHECK-NEXT: InitListExpr {{.*}} 'Accessor[2]'
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'Accessor'
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'Accessor'
-
-// Check __init functions are called
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}}__init
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}}__init
+// CHECK-NEXT:  VarDecl
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT:  VarDecl {{.*}} cinit
+// CHECK-NEXT:   MemberExpr
+// CHECK-NEXT:    DeclRefExpr
+// Init first accessor
+// CHECK-NEXT: CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .member_acc
+// CHECK-NEXT:        MemberExpr
+// CHECK-NEXT:         DeclRefExpr
+// CHECK-NEXT:     IntegerLiteral {{.*}} 0
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
+// CHECK-NEXT:  MemberExpr {{.*}}__init
+// Init second accessor
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .member_acc
+// CHECK-NEXT:        MemberExpr
+// CHECK-NEXT:         DeclRefExpr
+// CHECK-NEXT:     IntegerLiteral {{.*}} 1
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
+// CHECK-NEXT:  MemberExpr {{.*}}__init
 
 // Check Kernel_TemplatedStructArray parameters
 // CHECK: FunctionDecl {{.*}}Kernel_TemplatedStructArray{{.*}} 'void (S<int>)'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_s 'S<int>':'S<int>'
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
-// CHECK-NEXT: VarDecl {{.*}} used __SYCLKernel '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' cinit
-// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp{{.*}})'
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'S<int>':'S<int>' 'void (const S<int> &) noexcept'
+// CHECK-NEXT:  VarDecl
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT:  VarDecl {{.*}} cinit
+// CHECK-NEXT:   MemberExpr
+// CHECK-NEXT:    DeclRefExpr
+// CHECK-NEXT: CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} 'S<int>':'S<int>'
+// CHECK-NEXT:     DeclRefExpr
 // CHECK-NEXT: ImplicitCastExpr
-// CHECK-NEXT: DeclRefExpr {{.*}} 'S<int>':'S<int>' lvalue ParmVar {{.*}} '_arg_s' 'S<int>':'S<int>'
+// CHECK-NEXT:  UnaryOperator
+// CHECK-NEXT:    DeclRefExpr {{.*}} ParmVar {{.*}} '_arg_s'
+// CHECK-NEXT: IntegerLiteral {{.*}} 12
 
 // Check Kernel_Array_2D parameters
 // CHECK: FunctionDecl {{.*}}Kernel_Array_2D{{.*}} 'void (__wrapper_class)'
@@ -223,31 +253,23 @@ int main() {
 // Check Kernel_Array_2D inits
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
-// CHECK-NEXT: VarDecl {{.*}} cinit
-// CHECK-NEXT: InitListExpr
-// 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_array_2D' '__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_array_2D' '__wrapper_class'
-// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned
-// 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_array_2D' '__wrapper_class'
-// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned
-// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned
+// CHECK-NEXT:  VarDecl
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT:  VarDecl {{.*}} cinit
+// CHECK-NEXT:   MemberExpr
+// CHECK-NEXT:    DeclRefExpr
+// CHECK-NEXT: CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .array_2D
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT:  UnaryOperator
+// CHECK-NEXT:   MemberExpr
+// CHECK-NEXT:    DeclRefExpr {{.*}} ParmVar {{.*}} '_arg_array_2D'
+// CHECK-NEXT: IntegerLiteral {{.*}} 24
 
 // Check Kernel_NonDecomposedStruct parameters.
 // CHECK: FunctionDecl {{.*}}Kernel_NonDecomposedStruct{{.*}} 'void (__wrapper_class)'
@@ -255,20 +277,23 @@ int main() {
 // Check Kernel_NonDecomposedStruct inits
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
-// CHECK-NEXT: VarDecl {{.*}} cinit
-// CHECK-NEXT: InitListExpr
-// CHECK-NEXT: ArrayInitLoopExpr {{.*}} 'NonDecomposedStruct[2]'
-// CHECK-NEXT: OpaqueValueExpr {{.*}} 'NonDecomposedStruct[2]' lvalue
-// CHECK-NEXT: MemberExpr {{.*}} 'NonDecomposedStruct[2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_NonDecompStructArray' '__wrapper_class'
-// CHECK-NEXT: CXXConstructExpr {{.*}}'NonDecomposedStruct' 'void (const NonDecomposedStruct &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const NonDecomposedStruct' lvalue <NoOp>
-// CHECK-NEXT: ArraySubscriptExpr {{.*}}'NonDecomposedStruct' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'NonDecomposedStruct *' <ArrayToPointerDecay>
-// CHECK-NEXT: OpaqueValueExpr {{.*}} 'NonDecomposedStruct[2]' lvalue
-// CHECK-NEXT: MemberExpr {{.*}} 'NonDecomposedStruct[2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_NonDecompStructArray' '__wrapper_class'
-// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned
+// CHECK-NEXT:  VarDecl
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT:  VarDecl {{.*}} cinit
+// CHECK-NEXT:   MemberExpr
+// CHECK-NEXT:    DeclRefExpr
+// CHECK-NEXT: CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .NonDecompStructArray
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT:  UnaryOperator
+// CHECK-NEXT:   MemberExpr
+// CHECK-NEXT:    DeclRefExpr {{.*}} ParmVar {{.*}} '_arg_NonDecompStructArray'
+// CHECK-NEXT: IntegerLiteral {{.*}} 32
 
 // Check Kernel_StructWithPointers parameters.
 // CHECK: FunctionDecl {{.*}}Kernel_StructWithPointers{{.*}} 'void (__wrapper_class)'
@@ -276,152 +301,56 @@ int main() {
 // Check Kernel_StructWithPointers inits
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
-// CHECK-NEXT: VarDecl {{.*}} cinit
-// CHECK-NEXT: InitListExpr
-// CHECK-NEXT: InitListExpr {{.*}} 'StructWithPointers[2]'
-// Initializer for StructWithPointersArray[0]
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'StructWithPointers':'StructWithPointers' 'void (const StructWithPointers &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const StructWithPointers' lvalue <NoOp>
-// CHECK-NEXT: UnaryOperator {{.*}} 'StructWithPointers':'StructWithPointers' lvalue prefix '*' cannot overflow
-// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'StructWithPointers *' reinterpret_cast<StructWithPointers *> <BitCast>
-// CHECK-NEXT: UnaryOperator {{.*}} '__generated_StructWithPointers *' prefix '&' cannot overflow
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__generated_StructWithPointers' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__generated_StructWithPointers *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} '__generated_StructWithPointers[2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_StructWithPointersArray'
-// CHECK-NEXT: IntegerLiteral {{.*}} 0
-// Initializer for StructWithPointersArray[1]
-// CHECK: CXXConstructExpr {{.*}} 'StructWithPointers':'StructWithPointers' 'void (const StructWithPointers &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const StructWithPointers' lvalue <NoOp>
-// CHECK-NEXT: UnaryOperator {{.*}} 'StructWithPointers':'StructWithPointers' lvalue prefix '*' cannot overflow
-// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'StructWithPointers *' reinterpret_cast<StructWithPointers *> <BitCast>
-// CHECK-NEXT: UnaryOperator {{.*}} '__generated_StructWithPointers *' prefix '&' cannot overflow
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__generated_StructWithPointers' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__generated_StructWithPointers *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} '__generated_StructWithPointers[2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_StructWithPointersArray'
-// CHECK-NEXT: IntegerLiteral {{.*}} 1
+// CHECK-NEXT:  VarDecl
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT:  VarDecl {{.*}} cinit
+// CHECK-NEXT:   MemberExpr
+// CHECK-NEXT:    DeclRefExpr
+// CHECK-NEXT: CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .StructWithPointersArray
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT:  UnaryOperator
+// CHECK-NEXT:   MemberExpr
+// CHECK-NEXT:    DeclRefExpr {{.*}} ParmVar {{.*}} '_arg_StructWithPointersArray'
+// CHECK-NEXT: IntegerLiteral {{.*}} 48
 
 // Check Kernel_Array_Ptrs_2D parameters
 // CHECK: FunctionDecl {{.*}}Kernel_Array_Ptrs_2D{{.*}} 'void (__wrapper_class, __wrapper_class)'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPointers_2D '__wrapper_class'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPointers '__wrapper_class'
-
-// Check Kernel_Array_Ptrs_2D inits
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
-// CHECK-NEXT: VarDecl {{.*}} cinit
-// CHECK-NEXT: InitListExpr
-
-// Initializer for ArrayOfPointers_2D
-// CHECK-NEXT: InitListExpr {{.*}} 'int *[2][3]'
-// CHECK-NEXT: InitListExpr {{.*}} 'int *[3]'
-// Initializer for ArrayOfPointers_2D[0][0]
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
-// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
-// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
-// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D'
-// CHECK-NEXT: IntegerLiteral {{.*}} 0
-// CHECK-NEXT: IntegerLiteral {{.*}} 0
-
-// Initializer for ArrayOfPointers_2D[0][1]
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
-// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
-// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
-// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D'
-// CHECK-NEXT: IntegerLiteral {{.*}} 0
-// CHECK-NEXT: IntegerLiteral {{.*}} 1
-
-// Initializer for ArrayOfPointers_2D[0][2]
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
-// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
-// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
-// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D'
-// CHECK-NEXT: IntegerLiteral {{.*}} 0
-// CHECK-NEXT: IntegerLiteral {{.*}} 2
-
-// CHECK-NEXT: InitListExpr {{.*}} 'int *[3]'
-
-// Initializer for ArrayOfPointers_2D[1][0]
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
-// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
-// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
-// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D'
-// CHECK-NEXT: IntegerLiteral {{.*}} 1
-// CHECK-NEXT: IntegerLiteral {{.*}} 0
-
-// Initializer for ArrayOfPointers_2D[1][1]
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
-// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
-// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
-// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D'
-// CHECK-NEXT: IntegerLiteral {{.*}} 1
-// CHECK-NEXT: IntegerLiteral {{.*}} 1
-
-// Initializer for ArrayOfPointers_2D[1][2]
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
-// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
-// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
-// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *(*)[3]' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2][3]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers_2D'
-// CHECK-NEXT: IntegerLiteral {{.*}} 1
-// CHECK-NEXT: IntegerLiteral {{.*}} 2
-
-// Initializer for ArrayOfPointers
-// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]'
-// Initializer for ArrayOfPointers[0]
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
-// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
-// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
-// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers'
-// CHECK-NEXT: IntegerLiteral {{.*}} 0
-
-// Initializer for ArrayOfPointers[1]
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
-// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
-// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
-// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers'
-// CHECK-NEXT: IntegerLiteral {{.*}} 1
+// CHECK-NEXT:  VarDecl
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT:  VarDecl {{.*}} cinit
+// CHECK-NEXT:   MemberExpr
+// CHECK-NEXT:    DeclRefExpr
+// CHECK-NEXT: CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .ArrayOfPointers_2D
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT:  UnaryOperator
+// CHECK-NEXT:   MemberExpr
+// CHECK-NEXT:    DeclRefExpr {{.*}} ParmVar {{.*}} '_arg_ArrayOfPointers_2D'
+// CHECK-NEXT: IntegerLiteral {{.*}} 48
+// CHECK-NEXT: CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .ArrayOfPointers
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT:  UnaryOperator
+// CHECK-NEXT:   MemberExpr
+// CHECK-NEXT:    DeclRefExpr {{.*}} ParmVar {{.*}} '_arg_ArrayOfPointers'
+// CHECK-NEXT: IntegerLiteral {{.*}} 16
diff --git a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp
index 696590de8e5f6..f74a5aaa80205 100644
--- a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp
+++ b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp
@@ -34,7 +34,13 @@ int main() {
 // Check lambda declaration inside the wrapper
 
 // CHECK: DeclStmt
-// CHECK-NEXT: VarDecl {{.*}} used __SYCLKernel '(lambda at {{.*}}basic-kernel-wrapper.cpp{{.*}})'
+// CHECK: VarDecl {{.*}} '__wrapper_union'
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .readWriteAccessor
+// CHECK-NEXT:      DeclRefExpr
 
 // Check accessor initialization
 
diff --git a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp
index 0c3328559f898..3bdb706d69967 100644
--- a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp
+++ b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp
@@ -81,78 +81,114 @@ int main() {
 // CHECK: ParmVarDecl {{.*}} used _arg_some_const 'const int'
 
 // Check that lambda field of const built-in type is initialized
-// CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})'
-// CHECK-NEXT: InitListExpr
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'const int' lvalue ParmVar {{.*}} '_arg_some_const' 'const int'
+// CHECK:      VarDecl {{.*}} '__wrapper_union'
+// CHECK:      CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .some_const
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT:  UnaryOperator
+// CHECK-NEXT:    DeclRefExpr {{.*}} ParmVar {{.*}} '_arg_some_const'
+// CHECK-NEXT: IntegerLiteral {{.*}} 4
 
 // Check kernel parameters
 // CHECK: {{.*}}kernel_int{{.*}} 'void (int)'
 // CHECK: ParmVarDecl {{.*}} used _arg_data 'int'
 
 // Check that lambda field of built-in type is initialized
-// CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})'
-// CHECK-NEXT: InitListExpr
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_data' 'int'
+// CHECK: VarDecl {{.*}} '__wrapper_union'
+// CHECK:      CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .data
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT:  UnaryOperator
+// CHECK-NEXT:    DeclRefExpr {{.*}} ParmVar {{.*}} '_arg_data'
+// CHECK-NEXT: IntegerLiteral {{.*}} 4
 
 // Check kernel parameters
 // CHECK: {{.*}}kernel_struct{{.*}} 'void (__generated_test_struct)'
 // CHECK: ParmVarDecl {{.*}} used _arg_s '__generated_test_struct'
 
 // Check that lambda field of struct type is initialized
-// CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})'
-// CHECK-NEXT: InitListExpr {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})'
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'test_struct':'test_struct' 'void (const test_struct &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const test_struct' lvalue <NoOp>
-// CHECK-NEXT: UnaryOperator {{.*}} 'test_struct':'test_struct' lvalue prefix '*' cannot overflow
-// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'test_struct *' reinterpret_cast<test_struct *> <BitCast>
-// CHECK-NEXT: UnaryOperator {{.*}} '__generated_test_struct *' prefix '&' cannot overflow
-// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_test_struct' lvalue ParmVar {{.*}} '_arg_s'
+// CHECK: VarDecl {{.*}} '__wrapper_union'
+// CHECK:      CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .s
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT:  UnaryOperator
+// CHECK-NEXT:    DeclRefExpr {{.*}} ParmVar {{.*}} '_arg_s'
+// CHECK-NEXT: IntegerLiteral {{.*}} 80
 
 // Check kernel parameters
 // CHECK: {{.*}}kernel_pointer{{.*}} 'void (__global int *, __global int *, __wrapper_class)'
 // CHECK: ParmVarDecl {{.*}} used _arg_new_data_addr '__global int *'
 // CHECK: ParmVarDecl {{.*}} used _arg_data_addr '__global int *'
 // CHECK: ParmVarDecl {{.*}} used _arg_ptr_array '__wrapper_class'
-// CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})'
+// CHECK: VarDecl {{.*}} '__wrapper_union'
 
 // Check that lambda fields of pointer types are initialized
-// CHECK: InitListExpr
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_new_data_addr' '__global int *'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_data_addr' '__global int *'
-// CHECK: InitListExpr {{.*}} 'int *[2]'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
-// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
-// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
-// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array'
-// CHECK-NEXT: IntegerLiteral {{.*}} 0
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
-// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
-// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
-// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array'
-// CHECK-NEXT: IntegerLiteral {{.*}} 1
+// CHECK:      CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .new_data_addr
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion> 
+// CHECK-NEXT:     ImplicitCastExpr
+// CHECK-NEXT:      DeclRefExpr {{.*}} ParmVar {{.*}} '_arg_new_data_addr'
+// CHECK-NEXT:  IntegerLiteral {{.*}} 8
+// CHECK-NEXT: CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .data_addr
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion> 
+// CHECK-NEXT:     ImplicitCastExpr
+// CHECK-NEXT:      DeclRefExpr {{.*}} ParmVar {{.*}} '_arg_data_addr'
+// CHECK-NEXT:  IntegerLiteral {{.*}} 8
+// CHECK-NEXT: CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .ptr_array
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr
+// CHECK-NEXT:     DeclRefExpr {{.*}} ParmVar {{.*}} '_arg_ptr_array'
+// CHECK-NEXT:  IntegerLiteral {{.*}} 16
 
 // CHECK: FunctionDecl {{.*}}kernel_nns{{.*}} 'void (__generated_test_struct_simple)'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_tds '__generated_test_struct_simple'
 
-// CHECK: VarDecl {{.*}} used __SYCLKernel
-// CHECK: InitListExpr
-// CHECK: CXXConstructExpr {{.*}} 'Nested::TDS':'test_struct_simple' 'void (const test_struct_simple &) noexcept'
-// CHECK: ImplicitCastExpr {{.*}} 'const test_struct_simple' lvalue <NoOp>
-// CHECK: UnaryOperator {{.*}} 'Nested::TDS':'test_struct_simple' lvalue prefix '*' cannot overflow
-// CHECK: CXXReinterpretCastExpr {{.*}} 'Nested::TDS *' reinterpret_cast<struct Nested::TDS *> <BitCast>
-// CHECK: UnaryOperator {{.*}} '__generated_test_struct_simple *' prefix '&' cannot overflow
-// CHECK: DeclRefExpr {{.*}} '__generated_test_struct_simple' lvalue ParmVar {{.*}} '_arg_tds' '__generated_test_struct_simple'
+// CHECK: VarDecl {{.*}} __wrapper_union
+// CHECK:      CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .tds
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    DeclRefExpr {{.*}} ParmVar {{.*}} '_arg_tds'
+// CHECK-NEXT:  IntegerLiteral {{.*}} 16
diff --git a/clang/test/SemaSYCL/half-kernel-arg.cpp b/clang/test/SemaSYCL/half-kernel-arg.cpp
index 19577e5a578ca..cc1e987b8e9fc 100644
--- a/clang/test/SemaSYCL/half-kernel-arg.cpp
+++ b/clang/test/SemaSYCL/half-kernel-arg.cpp
@@ -20,8 +20,15 @@ int main() {
 // CHECK: {{.*}}kernel_half{{.*}} 'void (sycl::half)'
 // CHECK: ParmVarDecl {{.*}} used _arg_HostHalf 'sycl::half':'sycl::detail::half_impl::half'
 // // Check that lambda field of half type is initialized
-// CHECK: VarDecl {{.*}}'(lambda at {{.*}}'
-// CHECK-NEXT: InitListExpr {{.*}}'(lambda at {{.*}}'
-// CHECK-NEXT: CXXConstructExpr {{.*}}'sycl::detail::half_impl::half'{{.*}}
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::detail::half_impl::half'
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::half':'sycl::detail::half_impl::half' lvalue ParmVar {{.*}} '_arg_HostHalf' 'sycl::half':'sycl::detail::half_impl::half'
+// CHECK: VarDecl {{.*}} used __wrapper_union '__wrapper_union'
+// CHECK:      CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .HostHalf
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    DeclRefExpr {{.*}} '_arg_HostHalf'
+// CHECK-NEXT:  IntegerLiteral {{.*}} 2
diff --git a/clang/test/SemaSYCL/inheritance.cpp b/clang/test/SemaSYCL/inheritance.cpp
index 8776fe94a46ea..743762391fc7c 100644
--- a/clang/test/SemaSYCL/inheritance.cpp
+++ b/clang/test/SemaSYCL/inheritance.cpp
@@ -63,41 +63,77 @@ int main() {
 // Check initializers for derived and base classes.
 // Each class has it's own initializer list
 // Base classes should be initialized first.
-// CHECK: VarDecl {{.*}} used derived 'derived' cinit
-// CHECK-NEXT: InitListExpr {{.*}} 'derived'
+// CHECK: VarDecl {{.*}} __wrapper_union
 
 // base is a simple class with no corresponding generated type. Therefore
 // copy from ParamVar
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'base':'base' 'void (const base &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const base' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg__base' 'base'
+// CHECK:      CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    ImplicitCastExpr {{.*}} <DerivedToBase (base)>
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    DeclRefExpr {{.*}} '_arg__base'
+// CHECK-NEXT:  IntegerLiteral {{.*}} 12
 
 // second_base contains pointers and therefore the ParamVar is a new generated
 // type. Perform a copy of the corresponding kernel parameter via
 // reinterpret_cast.
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'second_base':'second_base' 'void (const second_base &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const second_base' lvalue <NoOp>
-// CHECK-NEXT: UnaryOperator {{.*}} 'second_base':'second_base' lvalue prefix '*' cannot overflow
-// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'second_base *' reinterpret_cast<second_base *> <BitCast>
-// CHECK-NEXT: UnaryOperator {{.*}} '__generated_second_base *' prefix '&' cannot overflow
-// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_second_base' lvalue ParmVar {{.*}} '_arg__base' '__generated_second_base'
+// CHECK-NEXT: CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    ImplicitCastExpr {{.*}} <DerivedToBase (second_base)>
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    DeclRefExpr {{.*}} '_arg__base' '__generated_second_base'
+// CHECK-NEXT:  IntegerLiteral {{.*}} 8
 
 // third_base contains special type accessor. Therefore it is decomposed and it's
 // data members are copied from corresponding ParamVar
-// CHECK-NEXT: InitListExpr {{.*}} 'third_base'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
-// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_d' '__wrapper_class'
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::accessor<char, 1, sycl::access::mode::read>'
-
-// Initialize fields of 'derived'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_a' 'int'
-
+// CHECK-NEXT: CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .d
+// CHECK-NEXT:     ImplicitCastExpr {{.*}} <DerivedToBase (third_base)>
+// CHECK-NEXT:      DeclRefExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    ImplicitCastExpr {{.*}} <AddressSpaceConversion>
+// CHECK-NEXT:     ImplicitCastExpr
+// CHECK-NEXT:      MemberExpr {{.*}} lvalue .
+// CHECK-NEXT:       DeclRefExpr {{.*}} '_arg_d' '__wrapper_class'
+// CHECK-NEXT:  IntegerLiteral {{.*}} 8
+// CHECK-NEXT: CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .AccField
+// CHECK-NEXT:     ImplicitCastExpr {{.*}} <DerivedToBase (third_base)>
+// CHECK-NEXT:      DeclRefExpr
 // Check kernel body for call to __init function of accessor
 // CHECK: CXXMemberCallExpr
 // CHECK-NEXT: MemberExpr {{.*}} lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} lvalue .AccField
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'third_base':'third_base' lvalue <DerivedToBase (third_base)>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'derived' lvalue Var {{.*}} 'derived' 'derived'
+// CHECK-NEXT: DeclRefExpr {{.*}} 'derived' lvalue Var {{.*}} '' 'derived &'
+
+// Initialize fields of 'derived'
+// CHECK:      CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .a
+// CHECK-NEXT:      DeclRefExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    DeclRefExpr {{.*}} '_arg_a'
+// CHECK-NEXT:  IntegerLiteral {{.*}} 4
diff --git a/clang/test/SemaSYCL/kernel-handler.cpp b/clang/test/SemaSYCL/kernel-handler.cpp
index 5a783dd69850d..94255dd90a1f4 100644
--- a/clang/test/SemaSYCL/kernel-handler.cpp
+++ b/clang/test/SemaSYCL/kernel-handler.cpp
@@ -35,10 +35,25 @@ int main() {
 // Check declaration and initialization of kernel object local clone
 // NONATIVESUPPORT-NEXT: CompoundStmt
 // NONATIVESUPPORT-NEXT: DeclStmt
-// NONATIVESUPPORT-NEXT: VarDecl {{.*}} cinit
-// NONATIVESUPPORT-NEXT: InitListExpr
-// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
-// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_a' 'int'
+// NONATIVESUPPORT-NEXT: VarDecl {{.*}}
+
+// Check declaration and initialization of reference to kernel object
+// NONATIVESUPPORT-NEXT: DeclStmt
+// NONATIVESUPPORT-NEXT: VarDecl {{.*}} '(lambda at {{.*}}) &' cinit 
+// NONATIVESUPPORT-NEXT: MemberExpr
+// NONATIVESUPPORT-NEXT: DeclRefExpr
+
+// NONATIVESUPPORT-NEXT: CallExpr
+// NONATIVESUPPORT-NEXT:  ImplicitCastExpr
+// NONATIVESUPPORT-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// NONATIVESUPPORT-NEXT:  ImplicitCastExpr
+// NONATIVESUPPORT-NEXT:   UnaryOperator
+// NONATIVESUPPORT-NEXT:    MemberExpr {{.*}} .a
+// NONATIVESUPPORT-NEXT:     DeclRefExpr
+// NONATIVESUPPORT-NEXT:  ImplicitCastExpr
+// NONATIVESUPPORT-NEXT:   UnaryOperator
+// NONATIVESUPPORT-NEXT:    DeclRefExpr {{.*}} '_arg_a'
+// NONATIVESUPPORT-NEXT:  IntegerLiteral {{.*}} 4
 
 // Check declaration and initialization of kernel handler local clone using default constructor
 // NONATIVESUPPORT-NEXT: DeclStmt
@@ -58,7 +73,7 @@ int main() {
 // NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'void (sycl::kernel_handler) const' lvalue CXXMethod {{.*}} 'operator()' 'void (sycl::kernel_handler) const'
 // Kernel body with clones
 // NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const (lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue
-// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})'
+// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}}) &'
 // NONATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'sycl::kernel_handler':'sycl::kernel_handler' 'void (const sycl::kernel_handler &) noexcept'
 // NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const sycl::kernel_handler' lvalue
 // NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}}'sycl::kernel_handler' lvalue Var {{.*}} 'kh' {{.*}}'sycl::kernel_handler'
@@ -71,12 +86,27 @@ int main() {
 // Check declaration and initialization of kernel object local clone
 // NONATIVESUPPORT-NEXT: CompoundStmt
 // NONATIVESUPPORT-NEXT: DeclStmt
-// NONATIVESUPPORT-NEXT: VarDecl {{.*}} cinit
-// NONATIVESUPPORT-NEXT: InitListExpr
-// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
-// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_a' 'int'
+// NONATIVESUPPORT-NEXT: VarDecl {{.*}}
 // NONATIVESUPPORT-NEXT: SYCLScopeAttr {{.*}} Implicit WorkGroup
 
+// Check declaration and initialization of reference to kernel object
+// NONATIVESUPPORT-NEXT: DeclStmt
+// NONATIVESUPPORT-NEXT: VarDecl {{.*}} '(lambda at {{.*}}) &' cinit 
+// NONATIVESUPPORT-NEXT: MemberExpr
+// NONATIVESUPPORT-NEXT: DeclRefExpr
+
+// NONATIVESUPPORT-NEXT: CallExpr
+// NONATIVESUPPORT-NEXT:  ImplicitCastExpr
+// NONATIVESUPPORT-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// NONATIVESUPPORT-NEXT:  ImplicitCastExpr
+// NONATIVESUPPORT-NEXT:   UnaryOperator
+// NONATIVESUPPORT-NEXT:    MemberExpr {{.*}} .a
+// NONATIVESUPPORT-NEXT:     DeclRefExpr
+// NONATIVESUPPORT-NEXT:  ImplicitCastExpr
+// NONATIVESUPPORT-NEXT:   UnaryOperator
+// NONATIVESUPPORT-NEXT:    DeclRefExpr {{.*}} '_arg_a'
+// NONATIVESUPPORT-NEXT:  IntegerLiteral {{.*}} 4
+
 // Check declaration and initialization of kernel handler local clone using default constructor
 // NONATIVESUPPORT-NEXT: DeclStmt
 // NONATIVESUPPORT-NEXT: VarDecl {{.*}} callinit
@@ -96,7 +126,7 @@ int main() {
 
 // Kernel body with clones
 // NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const (lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue
-// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})'
+// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}}) &'
 // NONATIVESUPPORT-NEXT: CXXTemporaryObjectExpr {{.*}} 'group<1>':'sycl::group<>' 'void () noexcept' zeroing
 // NONATIVESUPPORT-NEXT: CXXConstructExpr {{.*}}'sycl::kernel_handler' 'void (const sycl::kernel_handler &) noexcept'
 // NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}}'const sycl::kernel_handler' lvalue
@@ -112,10 +142,25 @@ int main() {
 // Check declaration and initialization of kernel object local clone
 // NATIVESUPPORT-NEXT: CompoundStmt
 // NATIVESUPPORT-NEXT: DeclStmt
-// NATIVESUPPORT-NEXT: VarDecl {{.*}} cinit
-// NATIVESUPPORT-NEXT: InitListExpr
-// NATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
-// NATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_a' 'int'
+// NATIVESUPPORT-NEXT: VarDecl {{.*}}
+
+// Check declaration and initialization of reference to kernel object
+// NATIVESUPPORT-NEXT: DeclStmt
+// NATIVESUPPORT-NEXT: VarDecl {{.*}} '(lambda at {{.*}}) &' cinit 
+// NATIVESUPPORT-NEXT: MemberExpr
+// NATIVESUPPORT-NEXT: DeclRefExpr
+
+// NATIVESUPPORT-NEXT: CallExpr
+// NATIVESUPPORT-NEXT:  ImplicitCastExpr
+// NATIVESUPPORT-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// NATIVESUPPORT-NEXT:  ImplicitCastExpr
+// NATIVESUPPORT-NEXT:   UnaryOperator
+// NATIVESUPPORT-NEXT:    MemberExpr {{.*}} .a
+// NATIVESUPPORT-NEXT:     DeclRefExpr
+// NATIVESUPPORT-NEXT:  ImplicitCastExpr
+// NATIVESUPPORT-NEXT:   UnaryOperator
+// NATIVESUPPORT-NEXT:    DeclRefExpr {{.*}} '_arg_a'
+// NATIVESUPPORT-NEXT:  IntegerLiteral {{.*}} 4
 
 // Check declaration and initialization of kernel handler local clone using default constructor
 // NATIVESUPPORT-NEXT: DeclStmt
@@ -127,7 +172,7 @@ int main() {
 
 // Kernel body with clones
 // NATIVESUPPORT: ImplicitCastExpr {{.*}} 'const (lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue
-// NATIVESUPPORT-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})'
+// NATIVESUPPORT-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}}) &'
 // NATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'sycl::kernel_handler':'sycl::kernel_handler' 'void (const sycl::kernel_handler &) noexcept'
 // NATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const sycl::kernel_handler' lvalue
 // NATIVESUPPORT-NEXT: DeclRefExpr {{.*}}'sycl::kernel_handler' lvalue Var {{.*}} 'kh' {{.*}}'sycl::kernel_handler'
diff --git a/clang/test/SemaSYCL/sampler.cpp b/clang/test/SemaSYCL/sampler.cpp
index 40506f800bf99..49979a14e8db5 100644
--- a/clang/test/SemaSYCL/sampler.cpp
+++ b/clang/test/SemaSYCL/sampler.cpp
@@ -29,7 +29,7 @@ int main() {
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__ocl_sampler_t)' lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::sampler':'sycl::sampler' lvalue
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}sampler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}sampler.cpp{{.*}})'
+// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}sampler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}sampler.cpp{{.*}}) &'
 //
 // Check the parameters of __init method
 // CHECK-NEXT: ImplicitCastExpr {{.*}} '__ocl_sampler_t':'sampler_t' <LValueToRValue>
diff --git a/clang/test/SemaSYCL/spec-const-kernel-arg.cpp b/clang/test/SemaSYCL/spec-const-kernel-arg.cpp
index a9ea7ba702ec3..35044b391e914 100644
--- a/clang/test/SemaSYCL/spec-const-kernel-arg.cpp
+++ b/clang/test/SemaSYCL/spec-const-kernel-arg.cpp
@@ -25,9 +25,24 @@ int main() {
 }
 
 // CHECK: FunctionDecl {{.*}}kernel_sc{{.*}} 'void ()'
-// CHECK: VarDecl {{.*}}'(lambda at {{.*}}'
-// CHECK-NEXT: InitListExpr {{.*}}'(lambda at {{.*}}'
-// CHECK-NEXT: CXXConstructExpr {{.*}}'sycl::ext::oneapi::experimental::spec_constant<char, class MyInt32Const>':'sycl::ext::oneapi::experimental::spec_constant<char, MyInt32Const>'
-// CHECK-NEXT: InitListExpr {{.*}} 'SpecConstantsWrapper'
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::ext::oneapi::experimental::spec_constant<int, class sc_name1>':'sycl::ext::oneapi::experimental::spec_constant<int, sc_name1>'
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::ext::oneapi::experimental::spec_constant<int, class sc_name2>':'sycl::ext::oneapi::experimental::spec_constant<int, sc_name2>'
+// CHECK: VarDecl {{.*}} used __wrapper_union '__wrapper_union'
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr {{.*}} 'sycl::ext::oneapi::experimental::spec_constant<char, MyInt32Const>'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .SC
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT: CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr {{.*}} 'sycl::ext::oneapi::experimental::spec_constant<int, sc_name1>'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .SC1
+// CHECK-NEXT:     MemberExpr {{.*}} .SCWrapper
+// CHECK-NEXT:      DeclRefExpr
+// CHECK-NEXT: CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr {{.*}} 'sycl::ext::oneapi::experimental::spec_constant<int, sc_name2>'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .SC2
+// CHECK-NEXT:     MemberExpr {{.*}} .SCWrapper
+// CHECK-NEXT:      DeclRefExpr
diff --git a/clang/test/SemaSYCL/spec-const-value-dependent-crash.cpp b/clang/test/SemaSYCL/spec-const-value-dependent-crash.cpp
index c44ac3fccaf1e..73b4b2895fa6c 100644
--- a/clang/test/SemaSYCL/spec-const-value-dependent-crash.cpp
+++ b/clang/test/SemaSYCL/spec-const-value-dependent-crash.cpp
@@ -20,5 +20,9 @@ int main() {
 
 // CHECK: FunctionDecl {{.*}}kernel_sc{{.*}} 'void ()'
 // CHECK: VarDecl {{.*}}'(lambda at {{.*}}'
-// CHECK-NEXT: InitListExpr {{.*}}'(lambda at {{.*}}'
-// CHECK-NEXT: CXXConstructExpr {{.*}}'sycl::ext::oneapi::experimental::spec_constant<int, class MyInt32Const>':'sycl::ext::oneapi::experimental::spec_constant<int, MyInt32Const>' 'void ()'
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr {{.*}} 'sycl::ext::oneapi::experimental::spec_constant<int, MyInt32Const>' 'void ()'
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     MemberExpr {{.*}} .SC
+// CHECK-NEXT:      DeclRefExpr
diff --git a/clang/test/SemaSYCL/stream.cpp b/clang/test/SemaSYCL/stream.cpp
index 1d220310c895d..3a69c59757ae0 100644
--- a/clang/test/SemaSYCL/stream.cpp
+++ b/clang/test/SemaSYCL/stream.cpp
@@ -50,146 +50,48 @@ int main() {
 
 // Initializers:
 
-// CHECK: InitListExpr {{.*}} '(lambda at
 // 'in_lambda'
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// 'in_lambda_array'
-// CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream[2]'
-// element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-
-// 'in_lambda_mdarray'
-// CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream[2][2]'
-// sub-array 0
-// CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream[2]'
-// element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// sub-array 1
-// CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream[2]'
-// element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-
-// HasStreams struct
-// CHECK: InitListExpr {{.*}} 'HasStreams'
-// HasStreams::s1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// HasStreams::s_array
-// CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream[2]'
-// element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-
-// HasArrayOfHasStreams
-// CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams'
-// HasArrayOfHasStreams::i
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar
-// HasArrayOfHasStreams::hs
-// CHECK-NEXT: InitListExpr {{.*}} 'HasStreams[2]'
-// HasStreams struct
-// CHECK-NEXT: InitListExpr {{.*}} 'HasStreams'
-// HasStreams::s1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// HasStreams::s_array
-// CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream[2]'
-// element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// HasStreams struct
-// CHECK-NEXT: InitListExpr {{.*}} 'HasStreams'
-// HasStreams::s1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// HasStreams::s_array
-// CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream[2]'
-// element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-
-// HasArrayOfHasStreams Array
-// CHECK: InitListExpr {{.*}} 'HasArrayOfHasStreams[2]'
-// // HasArrayOfHasStreams Struct
-// CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams'
-// HasArrayOfHasStreams::i
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar
-// HasArrayOfHasStreams::hs
-// CHECK-NEXT: InitListExpr {{.*}} 'HasStreams[2]'
-// HasStreams struct
-// CHECK-NEXT: InitListExpr {{.*}} 'HasStreams'
-// HasStreams::s1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// HasStreams::s_array
-// CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream[2]'
-// element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// HasStreams struct
-// CHECK-NEXT: InitListExpr {{.*}} 'HasStreams'
-// HasStreams::s1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// HasStreams::s_array
-// CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream[2]'
-// element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// HasArrayOfHasStreams Struct
-// CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams'
-// HasArrayOfHasStreams::i
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar
-// HasArrayOfHasStreams::hs
-// CHECK-NEXT: InitListExpr {{.*}} 'HasStreams[2]'
-// HasStreams struct
-// CHECK-NEXT: InitListExpr {{.*}} 'HasStreams'
-// HasStreams::s1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// HasStreams::s_array
-// CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream[2]'
-// element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// HasStreams struct
-// CHECK-NEXT: InitListExpr {{.*}} 'HasStreams'
-// HasStreams::s1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// HasStreams::s_array
-// CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream[2]'
-// element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-// element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
-
-// Calls to init
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     MemberExpr {{.*}} .in_lambda
+// CHECK-NEXT:      DeclRefExpr
 // in_lambda __init
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 
-// in_lambda_array
+// 'in_lambda_array'
 // element 0
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .in_lambda_array
+// CHECK-NEXT:        DeclRefExpr
+// CHECK-NEXT:     IntegerLiteral {{.*}} 0
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream[2]' lvalue .
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-
 // element 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .in_lambda_array
+// CHECK-NEXT:        DeclRefExpr
+// CHECK-NEXT:     IntegerLiteral {{.*}} 1
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
@@ -197,9 +99,21 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 
-// _in_lambda_mdarray
+// 'in_lambda_mdarray'
 // [0][0]
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       ArraySubscriptExpr
+// CHECK-NEXT:        ImplicitCastExpr
+// CHECK-NEXT:         MemberExpr {{.*}} .in_lambda_mdarray
+// CHECK-NEXT:          DeclRefExpr
+// CHECK-NEXT:       IntegerLiteral {{.*}} 0
+// CHECK-NEXT:     IntegerLiteral {{.*}} 0
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
@@ -209,8 +123,20 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// [0][1]
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// [1][0]
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       ArraySubscriptExpr
+// CHECK-NEXT:        ImplicitCastExpr
+// CHECK-NEXT:         MemberExpr {{.*}} .in_lambda_mdarray
+// CHECK-NEXT:          DeclRefExpr
+// CHECK-NEXT:       IntegerLiteral {{.*}} 0
+// CHECK-NEXT:     IntegerLiteral {{.*}} 1
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
@@ -220,8 +146,20 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// [1][0]
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// [0][1]
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       ArraySubscriptExpr
+// CHECK-NEXT:        ImplicitCastExpr
+// CHECK-NEXT:         MemberExpr {{.*}} .in_lambda_mdarray
+// CHECK-NEXT:          DeclRefExpr
+// CHECK-NEXT:       IntegerLiteral {{.*}} 1
+// CHECK-NEXT:     IntegerLiteral {{.*}} 0
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
@@ -232,7 +170,19 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // [1][1]
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       ArraySubscriptExpr
+// CHECK-NEXT:        ImplicitCastExpr
+// CHECK-NEXT:         MemberExpr {{.*}} .in_lambda_mdarray
+// CHECK-NEXT:          DeclRefExpr
+// CHECK-NEXT:       IntegerLiteral {{.*}} 1
+// CHECK-NEXT:     IntegerLiteral {{.*}} 1
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
@@ -243,14 +193,33 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 
-// HasStreams
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// HasStreams struct
+// HasStreams::s1
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     MemberExpr {{.*}} .s1
+// CHECK-NEXT:      MemberExpr {{.*}} .Struct
+// CHECK-NEXT:       DeclRefExpr
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
 // CHECK-NEXT: MemberExpr {{.*}}'HasStreams' lvalue .
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// array:
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// 'HasStreams::s_array'
+// element 0
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .s_array
+// CHECK-NEXT:        MemberExpr {{.*}} .Struct
+// CHECK-NEXT:         DeclRefExpr
+// CHECK-NEXT:     IntegerLiteral {{.*}} 0
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
@@ -259,7 +228,17 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // element 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .s_array
+// CHECK-NEXT:        MemberExpr {{.*}} .Struct
+// CHECK-NEXT:         DeclRefExpr
+// CHECK-NEXT:     IntegerLiteral {{.*}} 1
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
@@ -269,8 +248,34 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 
 // HasArrayOfHasStreams
-// First element
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// HasArrayOfHasStreams::i
+// CHECK:      CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .i
+// CHECK-NEXT:     MemberExpr {{.*}} .haohs
+// CHECK-NEXT:      DeclRefExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    DeclRefExpr {{.*}} '_arg_i'
+// CHECK-NEXT:  IntegerLiteral {{.*}} 4
+// HasArrayOfHasStreams::hs
+// HasStreams struct
+// HasStreams::s1
+// CHECK-NEXT: CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     MemberExpr {{.*}} .s1
+// CHECK-NEXT:      ArraySubscriptExpr
+// CHECK-NEXT:       ImplicitCastExpr
+// CHECK-NEXT:        MemberExpr {{.*}} .hs
+// CHECK-NEXT:         MemberExpr {{.*}} .haohs
+// CHECK-NEXT:          DeclRefExpr
+// CHECK-NEXT:      IntegerLiteral {{.*}} 0
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
 // CHECK-NEXT: ArraySubscriptExpr {{.*}}'HasStreams' lvalue
@@ -279,8 +284,23 @@ int main() {
 // CHECK-NEXT: MemberExpr {{.*}}'HasArrayOfHasStreams' lvalue .
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// array:
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// HasStreams::s_array
+// element 0
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .s_array
+// CHECK-NEXT:        ArraySubscriptExpr
+// CHECK-NEXT:         ImplicitCastExpr
+// CHECK-NEXT:          MemberExpr {{.*}} .hs
+// CHECK-NEXT:           MemberExpr {{.*}} .haohs
+// CHECK-NEXT:            DeclRefExpr
+// CHECK-NEXT:        IntegerLiteral {{.*}} 0
+// CHECK-NEXT:     IntegerLiteral {{.*}} 0
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
@@ -293,7 +313,21 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // element 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .s_array
+// CHECK-NEXT:        ArraySubscriptExpr
+// CHECK-NEXT:         ImplicitCastExpr
+// CHECK-NEXT:          MemberExpr {{.*}} .hs
+// CHECK-NEXT:           MemberExpr {{.*}} .haohs
+// CHECK-NEXT:            DeclRefExpr
+// CHECK-NEXT:        IntegerLiteral {{.*}} 0
+// CHECK-NEXT:     IntegerLiteral {{.*}} 1
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
@@ -305,8 +339,20 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// second element
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// HasStreams struct
+// HasStreams::s1
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     MemberExpr {{.*}} .s1
+// CHECK-NEXT:      ArraySubscriptExpr
+// CHECK-NEXT:       ImplicitCastExpr
+// CHECK-NEXT:        MemberExpr {{.*}} .hs
+// CHECK-NEXT:         MemberExpr {{.*}} .haohs
+// CHECK-NEXT:          DeclRefExpr
+// CHECK-NEXT:      IntegerLiteral {{.*}} 1
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
 // CHECK-NEXT: ArraySubscriptExpr {{.*}}'HasStreams' lvalue
@@ -315,8 +361,23 @@ int main() {
 // CHECK-NEXT: MemberExpr {{.*}}'HasArrayOfHasStreams' lvalue .
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// array:
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// HasStreams::s_array
+// element 0
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .s_array
+// CHECK-NEXT:        ArraySubscriptExpr
+// CHECK-NEXT:         ImplicitCastExpr
+// CHECK-NEXT:          MemberExpr {{.*}} .hs
+// CHECK-NEXT:           MemberExpr {{.*}} .haohs
+// CHECK-NEXT:            DeclRefExpr
+// CHECK-NEXT:        IntegerLiteral {{.*}} 1
+// CHECK-NEXT:     IntegerLiteral {{.*}} 0
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
@@ -329,7 +390,21 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // element 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .s_array
+// CHECK-NEXT:        ArraySubscriptExpr
+// CHECK-NEXT:         ImplicitCastExpr
+// CHECK-NEXT:          MemberExpr {{.*}} .hs
+// CHECK-NEXT:           MemberExpr {{.*}} .haohs
+// CHECK-NEXT:            DeclRefExpr
+// CHECK-NEXT:        IntegerLiteral {{.*}} 1
+// CHECK-NEXT:     IntegerLiteral {{.*}} 1
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
@@ -341,9 +416,43 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// HasArrayOfHasStreams array
-// First element
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+
+// HasArrayOfHasStreams Array
+// HasArrayOfHasStreams Struct
+// HasArrayOfHasStreams::i
+// CHECK:      CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .i
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .haohs
+// CHECK-NEXT:        DeclRefExpr
+// CHECK-NEXT:     IntegerLiteral {{.*}} 0
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    DeclRefExpr {{.*}} '_arg_i'
+// CHECK-NEXT:  IntegerLiteral {{.*}} 4
+// HasArrayOfHasStreams::hs
+// HasStreams struct
+// HasStreams::s1
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     MemberExpr {{.*}} .s1
+// CHECK-NEXT:      ArraySubscriptExpr
+// CHECK-NEXT:       ImplicitCastExpr
+// CHECK-NEXT:        MemberExpr {{.*}} .hs
+// CHECK-NEXT:         ArraySubscriptExpr
+// CHECK-NEXT:          ImplicitCastExpr
+// CHECK-NEXT:           MemberExpr {{.*}} .haohs_array
+// CHECK-NEXT:            DeclRefExpr
+// CHECK-NEXT:         IntegerLiteral {{.*}} 0
+// CHECK-NEXT:      IntegerLiteral {{.*}} 0
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
 // CHECK-NEXT: ArraySubscriptExpr {{.*}}'HasStreams' lvalue
@@ -355,8 +464,26 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// array:
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// HasStreams::s_array
+// element 0
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .s_array
+// CHECK-NEXT:        ArraySubscriptExpr
+// CHECK-NEXT:         ImplicitCastExpr
+// CHECK-NEXT:          MemberExpr {{.*}} .hs
+// CHECK-NEXT:           ArraySubscriptExpr
+// CHECK-NEXT:            ImplicitCastExpr
+// CHECK-NEXT:             MemberExpr {{.*}} .haohs_array
+// CHECK-NEXT:              DeclRefExpr
+// CHECK-NEXT:           IntegerLiteral {{.*}} 0
+// CHECK-NEXT:        IntegerLiteral {{.*}} 0
+// CHECK-NEXT:     IntegerLiteral {{.*}} 0
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
@@ -372,7 +499,24 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // element 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .s_array
+// CHECK-NEXT:        ArraySubscriptExpr
+// CHECK-NEXT:         ImplicitCastExpr
+// CHECK-NEXT:          MemberExpr {{.*}} .hs
+// CHECK-NEXT:           ArraySubscriptExpr
+// CHECK-NEXT:            ImplicitCastExpr
+// CHECK-NEXT:             MemberExpr {{.*}} .haohs_array
+// CHECK-NEXT:              DeclRefExpr
+// CHECK-NEXT:           IntegerLiteral {{.*}} 0
+// CHECK-NEXT:        IntegerLiteral {{.*}} 0
+// CHECK-NEXT:     IntegerLiteral {{.*}} 1
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
@@ -387,8 +531,23 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// second element
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// HasStreams struct
+// HasStreams::s1
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     MemberExpr {{.*}} .s1
+// CHECK-NEXT:      ArraySubscriptExpr
+// CHECK-NEXT:       ImplicitCastExpr
+// CHECK-NEXT:        MemberExpr {{.*}} .hs
+// CHECK-NEXT:         ArraySubscriptExpr
+// CHECK-NEXT:          ImplicitCastExpr
+// CHECK-NEXT:           MemberExpr {{.*}} .haohs_array
+// CHECK-NEXT:            DeclRefExpr
+// CHECK-NEXT:         IntegerLiteral {{.*}} 0
+// CHECK-NEXT:      IntegerLiteral {{.*}} 1
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
 // CHECK-NEXT: ArraySubscriptExpr {{.*}}'HasStreams' lvalue
@@ -400,8 +559,26 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// array:
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// HasStreams::s_array
+// element 0
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .s_array
+// CHECK-NEXT:        ArraySubscriptExpr
+// CHECK-NEXT:         ImplicitCastExpr
+// CHECK-NEXT:          MemberExpr {{.*}} .hs
+// CHECK-NEXT:           ArraySubscriptExpr
+// CHECK-NEXT:            ImplicitCastExpr
+// CHECK-NEXT:             MemberExpr {{.*}} .haohs_array
+// CHECK-NEXT:              DeclRefExpr
+// CHECK-NEXT:           IntegerLiteral {{.*}} 0
+// CHECK-NEXT:        IntegerLiteral {{.*}} 1
+// CHECK-NEXT:     IntegerLiteral {{.*}} 0
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
@@ -417,7 +594,24 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // element 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .s_array
+// CHECK-NEXT:        ArraySubscriptExpr
+// CHECK-NEXT:         ImplicitCastExpr
+// CHECK-NEXT:          MemberExpr {{.*}} .hs
+// CHECK-NEXT:           ArraySubscriptExpr
+// CHECK-NEXT:            ImplicitCastExpr
+// CHECK-NEXT:             MemberExpr {{.*}} .haohs_array
+// CHECK-NEXT:              DeclRefExpr
+// CHECK-NEXT:           IntegerLiteral {{.*}} 0
+// CHECK-NEXT:        IntegerLiteral {{.*}} 1
+// CHECK-NEXT:     IntegerLiteral {{.*}} 1
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}}  'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
@@ -432,8 +626,41 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// second element
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// HasArrayOfHasStreams Struct
+// HasArrayOfHasStreams::i
+// CHECK:      CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .i
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .haohs
+// CHECK-NEXT:        DeclRefExpr
+// CHECK-NEXT:     IntegerLiteral {{.*}} 1
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    DeclRefExpr {{.*}} '_arg_i'
+// CHECK-NEXT:  IntegerLiteral {{.*}} 4
+// HasArrayOfHasStreams::hs
+// HasStreams struct
+// HasStreams::s1
+// CHECK-NEXT: CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     MemberExpr {{.*}} .s1
+// CHECK-NEXT:      ArraySubscriptExpr
+// CHECK-NEXT:       ImplicitCastExpr
+// CHECK-NEXT:        MemberExpr {{.*}} .hs
+// CHECK-NEXT:         ArraySubscriptExpr
+// CHECK-NEXT:          ImplicitCastExpr
+// CHECK-NEXT:           MemberExpr {{.*}} .haohs_array
+// CHECK-NEXT:            DeclRefExpr
+// CHECK-NEXT:         IntegerLiteral {{.*}} 1
+// CHECK-NEXT:      IntegerLiteral {{.*}} 0
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}}  'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
 // CHECK-NEXT: ArraySubscriptExpr {{.*}}'HasStreams' lvalue
@@ -445,8 +672,26 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// array:
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// HasStreams::s_array
+// element 0
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .s_array
+// CHECK-NEXT:        ArraySubscriptExpr
+// CHECK-NEXT:         ImplicitCastExpr
+// CHECK-NEXT:          MemberExpr {{.*}} .hs
+// CHECK-NEXT:           ArraySubscriptExpr
+// CHECK-NEXT:            ImplicitCastExpr
+// CHECK-NEXT:             MemberExpr {{.*}} .haohs_array
+// CHECK-NEXT:              DeclRefExpr
+// CHECK-NEXT:           IntegerLiteral {{.*}} 1
+// CHECK-NEXT:        IntegerLiteral {{.*}} 0
+// CHECK-NEXT:     IntegerLiteral {{.*}} 0
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}}  'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
@@ -462,7 +707,24 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // element 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .s_array
+// CHECK-NEXT:        ArraySubscriptExpr
+// CHECK-NEXT:         ImplicitCastExpr
+// CHECK-NEXT:          MemberExpr {{.*}} .hs
+// CHECK-NEXT:           ArraySubscriptExpr
+// CHECK-NEXT:            ImplicitCastExpr
+// CHECK-NEXT:             MemberExpr {{.*}} .haohs_array
+// CHECK-NEXT:              DeclRefExpr
+// CHECK-NEXT:           IntegerLiteral {{.*}} 1
+// CHECK-NEXT:        IntegerLiteral {{.*}} 0
+// CHECK-NEXT:     IntegerLiteral {{.*}} 1
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}}  'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
@@ -477,7 +739,22 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// second element
+// HasStreams struct
+// HasStreams::s1
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     MemberExpr {{.*}} .s1
+// CHECK-NEXT:      ArraySubscriptExpr
+// CHECK-NEXT:       ImplicitCastExpr
+// CHECK-NEXT:        MemberExpr {{.*}} .hs
+// CHECK-NEXT:         ArraySubscriptExpr
+// CHECK-NEXT:          ImplicitCastExpr
+// CHECK-NEXT:           MemberExpr {{.*}} .haohs_array
+// CHECK-NEXT:            DeclRefExpr
+// CHECK-NEXT:         IntegerLiteral {{.*}} 1
+// CHECK-NEXT:      IntegerLiteral {{.*}} 1
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}}  'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
@@ -490,8 +767,26 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// array:
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// HasStreams::s_array
+// element 0
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .s_array
+// CHECK-NEXT:        ArraySubscriptExpr
+// CHECK-NEXT:         ImplicitCastExpr
+// CHECK-NEXT:          MemberExpr {{.*}} .hs
+// CHECK-NEXT:           ArraySubscriptExpr
+// CHECK-NEXT:            ImplicitCastExpr
+// CHECK-NEXT:             MemberExpr {{.*}} .haohs_array
+// CHECK-NEXT:              DeclRefExpr
+// CHECK-NEXT:           IntegerLiteral {{.*}} 1
+// CHECK-NEXT:        IntegerLiteral {{.*}} 1
+// CHECK-NEXT:     IntegerLiteral {{.*}} 0
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}}  'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
@@ -507,7 +802,24 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // element 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK:      CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     ArraySubscriptExpr
+// CHECK-NEXT:      ImplicitCastExpr
+// CHECK-NEXT:       MemberExpr {{.*}} .s_array
+// CHECK-NEXT:        ArraySubscriptExpr
+// CHECK-NEXT:         ImplicitCastExpr
+// CHECK-NEXT:          MemberExpr {{.*}} .hs
+// CHECK-NEXT:           ArraySubscriptExpr
+// CHECK-NEXT:            ImplicitCastExpr
+// CHECK-NEXT:             MemberExpr {{.*}} .haohs_array
+// CHECK-NEXT:              DeclRefExpr
+// CHECK-NEXT:           IntegerLiteral {{.*}} 1
+// CHECK-NEXT:        IntegerLiteral {{.*}} 1
+// CHECK-NEXT:     IntegerLiteral {{.*}} 1
+// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
 // CHECK-NEXT: MemberExpr {{.*}}  'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
@@ -523,6 +835,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 
+
 // Finalize
 // in_lambda __finalize
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
diff --git a/clang/test/SemaSYCL/union-kernel-param.cpp b/clang/test/SemaSYCL/union-kernel-param.cpp
index 312c0b8f05009..956d23afa04f2 100644
--- a/clang/test/SemaSYCL/union-kernel-param.cpp
+++ b/clang/test/SemaSYCL/union-kernel-param.cpp
@@ -30,8 +30,15 @@ int main() {
 // Check kernel inits
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
-// CHECK-NEXT: VarDecl {{.*}} cinit
-// CHECK-NEXT: InitListExpr
-// CHECK-NEXT: CXXConstructExpr {{.*}}'MyUnion' 'void (const MyUnion &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const MyUnion'
-// CHECK-NEXT: DeclRefExpr {{.*}}'MyUnion' lvalue ParmVar {{.*}} '_arg_accel' {{.*}}'MyUnion'
+// CHECK-NEXT: VarDecl {{.*}} __wrapper_union
+// CHECK:      CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .accel
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    DeclRefExpr {{.*}} '_arg_accel'
+// CHECK-NEXT:  IntegerLiteral {{.*}} 4
diff --git a/clang/test/SemaSYCL/union-kernel-param1.cpp b/clang/test/SemaSYCL/union-kernel-param1.cpp
index 48f5b0c7a7cb4..343940892179a 100644
--- a/clang/test/SemaSYCL/union-kernel-param1.cpp
+++ b/clang/test/SemaSYCL/union-kernel-param1.cpp
@@ -39,11 +39,18 @@ int main() {
 // Check kernel_A inits
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
-// CHECK-NEXT: VarDecl {{.*}} cinit
-// CHECK-NEXT: InitListExpr
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'union union_acc_t':'union_acc_t' 'void (const union_acc_t &) noexcept'
-// CHECK: ImplicitCastExpr {{.*}} 'const union_acc_t'
-// CHECK: DeclRefExpr {{.*}} 'union union_acc_t':'union_acc_t' lvalue ParmVar {{.*}} '_arg_union_acc' 'union union_acc_t':'union_acc_t'
+// CHECK-NEXT: VarDecl {{.*}} __wrapper_union
+// CHECK:      CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .union_acc
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    DeclRefExpr {{.*}} '_arg_union_acc'
+// CHECK-NEXT:  IntegerLiteral {{.*}} 8
 
 // Check kernel_B parameters
 // CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (S<int>)'
@@ -52,8 +59,15 @@ int main() {
 // Check kernel_B inits
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
-// CHECK-NEXT: VarDecl {{.*}} cinit
-// CHECK-NEXT: InitListExpr
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'S<int>':'S<int>' 'void (const S<int> &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}}  'const S<int>'
-// CHECK-NEXT: DeclRefExpr {{.*}} 'S<int>':'S<int>' lvalue ParmVar {{.*}} '_arg_s' 'S<int>':'S<int>'
+// CHECK-NEXT: VarDecl {{.*}} __wrapper_union
+// CHECK:      CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .s
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    DeclRefExpr {{.*}} '_arg_s'
+// CHECK-NEXT:  IntegerLiteral {{.*}} 12
diff --git a/clang/test/SemaSYCL/union-kernel-param2.cpp b/clang/test/SemaSYCL/union-kernel-param2.cpp
index 61abad21ddea8..041d13f546650 100644
--- a/clang/test/SemaSYCL/union-kernel-param2.cpp
+++ b/clang/test/SemaSYCL/union-kernel-param2.cpp
@@ -65,11 +65,18 @@ int main() {
 // Check kernel_A inits
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
-// CHECK-NEXT: VarDecl {{.*}} cinit
-// CHECK-NEXT: InitListExpr
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'union MyUnion':'MyUnion' 'void (const MyUnion &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const MyUnion'
-// CHECK-NEXT: DeclRefExpr {{.*}} 'union MyUnion':'MyUnion' lvalue ParmVar {{.*}} '_arg_union_mem' 'union MyUnion':'MyUnion'
+// CHECK-NEXT: VarDecl {{.*}} __wrapper_union
+// CHECK:      CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .union_mem
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    DeclRefExpr {{.*}} '_arg_union_mem'
+// CHECK-NEXT:  IntegerLiteral {{.*}} 20
 
 // Check kernel_B parameters
 // CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (union MyUnion, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
@@ -82,20 +89,33 @@ int main() {
 // Check kernel_B inits
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
-// CHECK-NEXT: VarDecl {{.*}} cinit
-// CHECK-NEXT: InitListExpr
-// CHECK-NEXT: InitListExpr {{.*}} 'MyStruct'
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'union MyUnion':'MyStruct::MyUnion' 'void (const MyStruct::MyUnion &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const MyStruct::MyUnion'
-// CHECK-NEXT: DeclRefExpr {{.*}} 'union MyUnion':'MyStruct::MyUnion' lvalue ParmVar {{.*}} '_arg_union_mem' 'union MyUnion':'MyStruct::MyUnion'
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::accessor<char, 1, sycl::access::mode::read>'
+// CHECK-NEXT: VarDecl {{.*}} __wrapper_union
+// CHECK:      CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .union_mem
+// CHECK-NEXT:     MemberExpr {{.*}} .struct_mem
+// CHECK-NEXT:      DeclRefExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    DeclRefExpr {{.*}} '_arg_union_mem'
+// CHECK-NEXT:  IntegerLiteral {{.*}} 12
+// CHECK-NEXT: CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:   ImplicitCastExpr
+// CHECK-NEXT:    UnaryOperator
+// CHECK-NEXT:     MemberExpr {{.*}} .AccField
+// CHECK-NEXT:      MemberExpr {{.*}} .struct_mem
+// CHECK-NEXT:       DeclRefExpr
 
 // Check call to __init to initialize AccField
 // CHECK-NEXT: CXXMemberCallExpr
 // CHECK-NEXT: MemberExpr {{.*}} lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} lvalue .AccField
 // CHECK-NEXT: MemberExpr {{.*}} lvalue .struct_mem
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}union-kernel-param2.cpp:48:9)' lvalue Var {{.*}} '__SYCLKernel' '(lambda at {{.*}}union-kernel-param2.cpp:48:9)'
+// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}union-kernel-param2.cpp:48:9)' lvalue Var {{.*}} '' '(lambda at {{.*}}union-kernel-param2.cpp:48:9) &'
 
 // Check kernel_C parameters
 // CHECK: FunctionDecl {{.*}}kernel_C{{.*}} 'void (__generated_MyStructWithPtr)'
@@ -104,11 +124,15 @@ int main() {
 // Check kernel_C inits
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
-// CHECK-NEXT: VarDecl {{.*}} cinit
-// CHECK-NEXT: InitListExpr
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'struct MyStructWithPtr':'MyStructWithPtr' 'void (const MyStructWithPtr &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const MyStructWithPtr' lvalue <NoOp>
-// CHECK-NEXT: UnaryOperator {{.*}} 'struct MyStructWithPtr':'MyStructWithPtr' lvalue prefix '*' cannot overflow
-// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'struct MyStructWithPtr *' reinterpret_cast<struct MyStructWithPtr *> <BitCast>
-// CHECK-NEXT: UnaryOperator {{.*}} '__generated_MyStructWithPtr *' prefix '&' cannot overflow
-// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_MyStructWithPtr' lvalue ParmVar {{.*}} '_arg_structWithPtr_mem' '__generated_MyStructWithPtr'
+// CHECK-NEXT: VarDecl {{.*}} __wrapper_union
+// CHECK:      CallExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   DeclRefExpr {{.*}} '__builtin_memcpy'
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .structWithPtr_mem
+// CHECK-NEXT:     DeclRefExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    DeclRefExpr {{.*}} '_arg_structWithPtr_mem'
+// CHECK-NEXT:  IntegerLiteral {{.*}} 24
diff --git a/clang/test/SemaSYCL/wrapped-accessor.cpp b/clang/test/SemaSYCL/wrapped-accessor.cpp
index 5d633dbd454aa..b2ab45a18f64d 100644
--- a/clang/test/SemaSYCL/wrapped-accessor.cpp
+++ b/clang/test/SemaSYCL/wrapped-accessor.cpp
@@ -35,17 +35,28 @@ int main() {
 
 // Check that wrapper object itself is initialized with corresponding kernel
 // argument
-// CHECK: VarDecl {{.*}}'(lambda at {{.*}}wrapped-accessor.cpp{{.*}})'
-// CHECK-NEXT: InitListExpr {{.*}}'(lambda at {{.*}}wrapped-accessor.cpp{{.*}})'
-// CHECK-NEXT: InitListExpr {{.*}}'AccWrapper<sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>>'
-// CHECK-NEXT: CXXConstructExpr {{.*}}'sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>':'sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>' 'void () noexcept'
+// CHECK: VarDecl {{.*}} '__wrapper_union'
+
+// Check that the ref is generated properly
+// CHECK: VarDecl {{.*}} used '(lambda at {{.*}}wrapped-accessor.cpp{{.*}}) &'
+// CHECK-NEXT: MemberExpr {{.*}} '(lambda at {{.*}}wrapped-accessor.cpp{{.*}})' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_union' lvalue Var {{.*}} '__wrapper_union' '__wrapper_union'
+
+// Build accessor
+// CHECK-NEXT: CXXNewExpr
+// CHECK-NEXT:  CXXConstructExpr
+// CHECK-NEXT:  ImplicitCastExpr
+// CHECK-NEXT:   UnaryOperator
+// CHECK-NEXT:    MemberExpr {{.*}} .accessor
+// CHECK-NEXT:     MemberExpr {{.*}} .acc_wrapped
+// CHECK-NEXT:      DeclRefExpr
 
 // Check that accessor field of the wrapper object is initialized using __init method
 // CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ({{.*}}PtrType, range<1>, range<1>, id<1>)' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>':'sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>' lvalue .accessor {{.*}}
-// CHECK-NEXT: MemberExpr {{.*}} 'AccWrapper<decltype(acc)>':'AccWrapper<sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>>' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}wrapped-accessor.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}wrapped-accessor.cpp{{.*}})'
+// CHECK-NEXT:  MemberExpr {{.*}} 'void ({{.*}}PtrType, range<1>, range<1>, id<1>)' lvalue .__init
+// CHECK-NEXT:   MemberExpr {{.*}} 'sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>':'sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>' lvalue .accessor {{.*}}
+// CHECK-NEXT:    MemberExpr {{.*}} 'AccWrapper<decltype(acc)>':'AccWrapper<sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>>' lvalue .
+// CHECK-NEXT:     DeclRefExpr {{.*}} '(lambda at {{.*}}wrapped-accessor.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}wrapped-accessor.cpp{{.*}}) &'
 
 // Parameters of the _init method
 // CHECK-NEXT: ImplicitCastExpr {{.*}} <LValueToRValue>