diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 3b3976a21fc01..7f85562f94822 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -76,6 +76,10 @@ class Util {
   /// stream class.
   static bool isSyclStreamType(const QualType &Ty);
 
+  /// Checks whether given clang type is a full specialization of the SYCL
+  /// half class.
+  static bool isSyclHalfType(const QualType &Ty);
+
   /// Checks whether given clang type is a standard SYCL API class with given
   /// name.
   /// \param Ty    the clang type being checked
@@ -728,9 +732,6 @@ constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc,
 // anonymous namespace so these don't get linkage.
 namespace {
 
-QualType getItemType(const FieldDecl *FD) { return FD->getType(); }
-QualType getItemType(const CXXBaseSpecifier &BS) { return BS.getType(); }
-
 // These enable handler execution only when previous handlers succeed.
 template <typename... Tn>
 static bool handleField(FieldDecl *FD, QualType FDTy, Tn &&... tn) {
@@ -774,11 +775,6 @@ template <typename T> using bind_param_t = typename bind_param<T>::type;
 //              })...)
 
 // Implements the 'for-each-visitor'  pattern.
-template <typename ParentTy, typename... Handlers>
-static void VisitAccessorWrapper(CXXRecordDecl *Owner, ParentTy &Parent,
-                                 CXXRecordDecl *Wrapper,
-                                 Handlers &... handlers);
-
 template <typename RangeTy, typename... Handlers>
 static void VisitField(CXXRecordDecl *Owner, RangeTy &&Item, QualType ItemTy,
                        Handlers &... handlers) {
@@ -788,26 +784,15 @@ static void VisitField(CXXRecordDecl *Owner, RangeTy &&Item, QualType ItemTy,
     KF_FOR_EACH(handleSyclStreamType, Item, ItemTy);
   else if (Util::isSyclSamplerType(ItemTy))
     KF_FOR_EACH(handleSyclSamplerType, Item, ItemTy);
+  else if (Util::isSyclHalfType(ItemTy))
+    KF_FOR_EACH(handleSyclHalfType, Item, ItemTy);
   else if (ItemTy->isStructureOrClassType())
-    VisitAccessorWrapper(Owner, Item, ItemTy->getAsCXXRecordDecl(),
-                         handlers...);
-    // FIXME Enable this when structs are replaced by their fields
-#define STRUCTS_DECOMPOSED 0
-#if STRUCTS_DECOMPOSED
+    VisitRecord(Owner, Item, ItemTy->getAsCXXRecordDecl(), handlers...);
   else if (ItemTy->isArrayType())
     VisitArrayElements(Item, ItemTy, handlers...);
   else if (ItemTy->isScalarType())
     KF_FOR_EACH(handleScalarType, Item, ItemTy);
 }
-#else
-}
-
-template <typename RangeTy, typename... Handlers>
-static void VisitScalarField(CXXRecordDecl *Owner, RangeTy &&Item,
-                             QualType ItemTy, Handlers &... handlers) {
-  KF_FOR_EACH(handleScalarType, Item, ItemTy);
-}
-#endif
 
 template <typename RangeTy, typename... Handlers>
 static void VisitArrayElements(RangeTy Item, QualType FieldTy,
@@ -817,71 +802,109 @@ static void VisitArrayElements(RangeTy Item, QualType FieldTy,
   int64_t ElemCount = CAT->getSize().getSExtValue();
   std::initializer_list<int>{(handlers.enterArray(), 0)...};
   for (int64_t Count = 0; Count < ElemCount; Count++) {
-#if STRUCTS_DECOMPOSED
     VisitField(nullptr, Item, ET, handlers...);
-#else
-    if (ET->isScalarType())
-      VisitScalarField(nullptr, Item, ET, handlers...);
-    else
-      VisitField(nullptr, Item, ET, handlers...);
-#endif
     (void)std::initializer_list<int>{(handlers.nextElement(ET), 0)...};
   }
   (void)std::initializer_list<int>{
       (handlers.leaveArray(Item, ET, ElemCount), 0)...};
 }
 
-template <typename RangeTy, typename... Handlers>
-static void VisitAccessorWrapperHelper(CXXRecordDecl *Owner, RangeTy Range,
-                                       Handlers &... handlers) {
-  for (const auto &Item : Range) {
-    QualType ItemTy = getItemType(Item);
-    (void)std::initializer_list<int>{(handlers.enterField(Owner, Item), 0)...};
-    VisitField(Owner, Item, ItemTy, handlers...);
-    (void)std::initializer_list<int>{(handlers.leaveField(Owner, Item), 0)...};
+template <typename ParentTy, typename... Handlers>
+static void VisitRecord(CXXRecordDecl *Owner, ParentTy &Parent,
+                        CXXRecordDecl *Wrapper, Handlers &... handlers);
+
+template <typename... Handlers>
+static void VisitRecordHelper(CXXRecordDecl *Owner,
+                              clang::CXXRecordDecl::base_class_range Range,
+                              Handlers &... handlers) {
+  for (const auto &Base : Range) {
+    (void)std::initializer_list<int>{(handlers.enterField(Owner, Base), 0)...};
+    QualType BaseTy = Base.getType();
+    // Handle accessor class as base
+    if (Util::isSyclAccessorType(BaseTy)) {
+      (void)std::initializer_list<int>{
+          (handlers.handleSyclAccessorType(Base, BaseTy), 0)...};
+    } else if (Util::isSyclStreamType(BaseTy)) {
+      // Handle stream class as base
+      (void)std::initializer_list<int>{
+          (handlers.handleSyclStreamType(Base, BaseTy), 0)...};
+    } else
+      // For all other bases, visit the record
+      VisitRecord(Owner, Base, BaseTy->getAsCXXRecordDecl(), handlers...);
+    (void)std::initializer_list<int>{(handlers.leaveField(Owner, Base), 0)...};
   }
 }
 
+template <typename... Handlers>
+static void VisitRecordHelper(CXXRecordDecl *Owner,
+                              clang::RecordDecl::field_range Range,
+                              Handlers &... handlers) {
+  VisitRecordFields(Owner, handlers...);
+}
+
 // Parent contains the FieldDecl or CXXBaseSpecifier that was used to enter
 // the Wrapper structure that we're currently visiting. Owner is the parent
 // type (which doesn't exist in cases where it is a FieldDecl in the
 // 'root'), and Wrapper is the current struct being unwrapped.
 template <typename ParentTy, typename... Handlers>
-static void VisitAccessorWrapper(CXXRecordDecl *Owner, ParentTy &Parent,
-                                 CXXRecordDecl *Wrapper,
-                                 Handlers &... handlers) {
+static void VisitRecord(CXXRecordDecl *Owner, ParentTy &Parent,
+                        CXXRecordDecl *Wrapper, Handlers &... handlers) {
   (void)std::initializer_list<int>{(handlers.enterStruct(Owner, Parent), 0)...};
-  VisitAccessorWrapperHelper(Wrapper, Wrapper->bases(), handlers...);
-  VisitAccessorWrapperHelper(Wrapper, Wrapper->fields(), handlers...);
+  VisitRecordHelper(Wrapper, Wrapper->bases(), handlers...);
+  VisitRecordHelper(Wrapper, Wrapper->fields(), handlers...);
+  (void)std::initializer_list<int>{(handlers.leaveStruct(Owner, Parent), 0)...};
+}
+
+// FIXME: Can this be refactored/handled some other way?
+template <typename ParentTy, typename... Handlers>
+static void VisitStreamRecord(CXXRecordDecl *Owner, ParentTy &Parent,
+                              CXXRecordDecl *Wrapper, Handlers &... handlers) {
+  (void)std::initializer_list<int>{(handlers.enterStruct(Owner, Parent), 0)...};
+  for (const auto &Field : Wrapper->fields()) {
+    QualType FieldTy = Field->getType();
+    (void)std::initializer_list<int>{
+        (handlers.enterField(Wrapper, Field), 0)...};
+    // Required to initialize accessors inside streams.
+    if (Util::isSyclAccessorType(FieldTy))
+      KF_FOR_EACH(handleSyclAccessorType, Field, FieldTy);
+    (void)std::initializer_list<int>{
+        (handlers.leaveField(Wrapper, Field), 0)...};
+  }
   (void)std::initializer_list<int>{(handlers.leaveStruct(Owner, Parent), 0)...};
 }
 
+template <typename... Handlers>
+static void VisitRecordBases(CXXRecordDecl *KernelFunctor,
+                             Handlers &... handlers) {
+  VisitRecordHelper(KernelFunctor, KernelFunctor->bases(), handlers...);
+}
+
 // A visitor function that dispatches to functions as defined in
 // SyclKernelFieldHandler for the purposes of kernel generation.
 template <typename... Handlers>
-static void VisitRecordFields(RecordDecl::field_range Fields,
-                              Handlers &... handlers) {
+static void VisitRecordFields(CXXRecordDecl *Owner, Handlers &... handlers) {
 
-  for (const auto Field : Fields) {
-    (void)std::initializer_list<int>{
-        (handlers.enterField(nullptr, Field), 0)...};
+  for (const auto Field : Owner->fields()) {
+    (void)std::initializer_list<int>{(handlers.enterField(Owner, Field), 0)...};
     QualType FieldTy = Field->getType();
 
     if (Util::isSyclAccessorType(FieldTy))
       KF_FOR_EACH(handleSyclAccessorType, Field, FieldTy);
     else if (Util::isSyclSamplerType(FieldTy))
       KF_FOR_EACH(handleSyclSamplerType, Field, FieldTy);
+    else if (Util::isSyclHalfType(FieldTy))
+      KF_FOR_EACH(handleSyclHalfType, Field, FieldTy);
     else if (Util::isSyclSpecConstantType(FieldTy))
       KF_FOR_EACH(handleSyclSpecConstantType, Field, FieldTy);
     else if (Util::isSyclStreamType(FieldTy)) {
-      // Stream actually wraps accessors, so do recursion
       CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl();
-      VisitAccessorWrapper(nullptr, Field, RD, handlers...);
+      // Handle accessors in stream class.
+      VisitStreamRecord(Owner, Field, RD, handlers...);
       KF_FOR_EACH(handleSyclStreamType, Field, FieldTy);
     } else if (FieldTy->isStructureOrClassType()) {
       if (KF_FOR_EACH(handleStructType, Field, FieldTy)) {
         CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl();
-        VisitAccessorWrapper(nullptr, Field, RD, handlers...);
+        VisitRecord(Owner, Field, RD, handlers...);
       }
     } else if (FieldTy->isReferenceType())
       KF_FOR_EACH(handleReferenceType, Field, FieldTy);
@@ -890,12 +913,11 @@ static void VisitRecordFields(RecordDecl::field_range Fields,
     else if (FieldTy->isArrayType()) {
       if (KF_FOR_EACH(handleArrayType, Field, FieldTy))
         VisitArrayElements(Field, FieldTy, handlers...);
-    } else if (FieldTy->isScalarType())
+    } else if (FieldTy->isScalarType() || FieldTy->isVectorType())
       KF_FOR_EACH(handleScalarType, Field, FieldTy);
     else
       KF_FOR_EACH(handleOtherType, Field, FieldTy);
-    (void)std::initializer_list<int>{
-        (handlers.leaveField(nullptr, Field), 0)...};
+    (void)std::initializer_list<int>{(handlers.leaveField(Owner, Field), 0)...};
   }
 #undef KF_FOR_EACH
 } // namespace
@@ -927,6 +949,10 @@ template <typename Derived> class SyclKernelFieldHandler {
     return true;
   }
   virtual bool handleSyclStreamType(FieldDecl *, QualType) { return true; }
+  virtual bool handleSyclHalfType(const CXXBaseSpecifier &, QualType) {
+    return true;
+  }
+  virtual bool handleSyclHalfType(FieldDecl *, QualType) { return true; }
   virtual bool handleStructType(FieldDecl *, QualType) { return true; }
   virtual bool handleReferenceType(FieldDecl *, QualType) { return true; }
   virtual bool handlePointerType(FieldDecl *, QualType) { return true; }
@@ -1223,8 +1249,7 @@ class SyclKernelDeclCreator
     return true;
   }
 
-  // FIXME Remove this function when structs are replaced by their fields
-  bool handleStructType(FieldDecl *FD, QualType FieldTy) final {
+  bool handleSyclHalfType(FieldDecl *FD, QualType FieldTy) final {
     addParam(FD, FieldTy);
     return true;
   }
@@ -1248,8 +1273,8 @@ class SyclKernelDeclCreator
     return ArrayRef<ParmVarDecl *>(std::begin(Params) + LastParamIndex,
                                    std::end(Params));
   }
-
   using SyclKernelFieldHandler::handleScalarType;
+  using SyclKernelFieldHandler::handleSyclHalfType;
   using SyclKernelFieldHandler::handleSyclSamplerType;
 };
 
@@ -1440,24 +1465,40 @@ class SyclKernelBodyCreator
 
   bool handleSpecialType(FieldDecl *FD, QualType Ty) {
     const auto *RecordDecl = Ty->getAsCXXRecordDecl();
-    ArraySubscriptExpr *ArrayRef =
-        dyn_cast<ArraySubscriptExpr>(MemberExprBases.back());
-    // Perform initialization only if decomposed from array
-    if (ArrayRef || MemberExprBases.size() == 2) {
-      InitializedEntity Entity =
-          InitializedEntity::InitializeMember(FD, &VarEntity);
-      // Initialize with the default constructor.
-      InitializationKind InitKind =
-          InitializationKind::CreateDefault(SourceLocation());
-      InitializationSequence InitSeq(SemaRef, Entity, InitKind, None);
-      ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, None);
-      InitExprs.push_back(MemberInit.get());
-    }
+    // TODO: VarEntity is initialized entity for KernelObjClone, I guess we need
+    // to create new one when enter new struct.
+    InitializedEntity Entity =
+        InitializedEntity::InitializeMember(FD, &VarEntity);
+    // Initialize with the default constructor.
+    InitializationKind InitKind =
+        InitializationKind::CreateDefault(SourceLocation());
+    InitializationSequence InitSeq(SemaRef, Entity, InitKind, None);
+    ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, None);
+    InitExprs.push_back(MemberInit.get());
+
     createSpecialMethodCall(RecordDecl, MemberExprBases.back(), InitMethodName,
                             FD);
     return true;
   }
 
+  bool handleSpecialType(const CXXBaseSpecifier &BS, QualType Ty) {
+    const auto *RecordDecl = Ty->getAsCXXRecordDecl();
+    // TODO: VarEntity is initialized entity for KernelObjClone, I guess we need
+    // to create new one when enter new struct.
+    InitializedEntity Entity = InitializedEntity::InitializeBase(
+        SemaRef.Context, &BS, /*IsInheritedVirtualBase*/ false, &VarEntity);
+    // Initialize with the default constructor.
+    InitializationKind InitKind =
+        InitializationKind::CreateDefault(SourceLocation());
+    InitializationSequence InitSeq(SemaRef, Entity, InitKind, None);
+    ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, None);
+    InitExprs.push_back(MemberInit.get());
+
+    createSpecialMethodCall(RecordDecl, MemberExprBases.back(), InitMethodName,
+                            nullptr);
+    return true;
+  }
+
 public:
   SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC,
                         CXXRecordDecl *KernelObj,
@@ -1489,9 +1530,7 @@ class SyclKernelBodyCreator
   }
 
   bool handleSyclAccessorType(const CXXBaseSpecifier &BS, QualType Ty) final {
-    // FIXME SYCL accessor should be usable as a base type
-    // See https://github.com/intel/llvm/issues/28.
-    return true;
+    return handleSpecialType(BS, Ty);
   }
 
   bool handleSyclSamplerType(FieldDecl *FD, QualType Ty) final {
@@ -1515,13 +1554,12 @@ class SyclKernelBodyCreator
     return true;
   }
 
-  bool handlePointerType(FieldDecl *FD, QualType FieldTy) final {
+  bool handleSyclHalfType(FieldDecl *FD, QualType Ty) final {
     createExprForStructOrScalar(FD);
     return true;
   }
 
-  // FIXME Remove this function when structs are replaced by their fields
-  bool handleStructType(FieldDecl *FD, QualType FieldTy) final {
+  bool handlePointerType(FieldDecl *FD, QualType FieldTy) final {
     createExprForStructOrScalar(FD);
     return true;
   }
@@ -1534,6 +1572,68 @@ class SyclKernelBodyCreator
     return true;
   }
 
+  bool enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final {
+    CXXCastPath BasePath;
+    QualType DerivedTy(RD->getTypeForDecl(), 0);
+    QualType BaseTy = BS.getType();
+    SemaRef.CheckDerivedToBaseConversion(DerivedTy, BaseTy, SourceLocation(),
+                                         SourceRange(), &BasePath,
+                                         /*IgnoreBaseAccess*/ true);
+    auto Cast = ImplicitCastExpr::Create(
+        SemaRef.Context, BaseTy, CK_DerivedToBase, MemberExprBases.back(),
+        /* CXXCastPath=*/&BasePath, VK_LValue);
+    MemberExprBases.push_back(Cast);
+    return true;
+  }
+
+  void addStructInit(const CXXRecordDecl *RD) {
+    if (!RD)
+      return;
+
+    const ASTRecordLayout &Info =
+        SemaRef.getASTContext().getASTRecordLayout(RD);
+    int NumberOfFields = Info.getFieldCount();
+    int popOut = NumberOfFields + RD->getNumBases();
+
+    llvm::SmallVector<Expr *, 16> BaseInitExprs;
+    for (int I = 0; I < popOut; I++) {
+      BaseInitExprs.push_back(InitExprs.back());
+      InitExprs.pop_back();
+    }
+    std::reverse(BaseInitExprs.begin(), BaseInitExprs.end());
+
+    Expr *ILE = new (SemaRef.getASTContext())
+        InitListExpr(SemaRef.getASTContext(), SourceLocation(), BaseInitExprs,
+                     SourceLocation());
+    ILE->setType(QualType(RD->getTypeForDecl(), 0));
+    InitExprs.push_back(ILE);
+  }
+
+  bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD) final {
+    const CXXRecordDecl *RD = FD->getType()->getAsCXXRecordDecl();
+
+    // Initializers for accessors inside stream not added.
+    if (!Util::isSyclStreamType(FD->getType()))
+      addStructInit(RD);
+    // Pop out unused initializers created in handleSyclAccesorType
+    // for accessors inside stream class.
+    else {
+      for (const auto &Field : RD->fields()) {
+        QualType FieldTy = Field->getType();
+        if (Util::isSyclAccessorType(FieldTy))
+          InitExprs.pop_back();
+      }
+    }
+    return true;
+  }
+
+  bool leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final {
+    const CXXRecordDecl *BaseClass = BS.getType()->getAsCXXRecordDecl();
+    addStructInit(BaseClass);
+    MemberExprBases.pop_back();
+    return true;
+  }
+
   bool enterField(const CXXRecordDecl *RD, FieldDecl *FD) final {
     if (!FD->getType()->isReferenceType())
       MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD));
@@ -1580,10 +1680,13 @@ class SyclKernelBodyCreator
 
   using SyclKernelFieldHandler::enterArray;
   using SyclKernelFieldHandler::enterField;
+  using SyclKernelFieldHandler::enterStruct;
   using SyclKernelFieldHandler::handleScalarType;
+  using SyclKernelFieldHandler::handleSyclHalfType;
   using SyclKernelFieldHandler::handleSyclSamplerType;
   using SyclKernelFieldHandler::leaveArray;
   using SyclKernelFieldHandler::leaveField;
+  using SyclKernelFieldHandler::leaveStruct;
 };
 
 class SyclKernelIntHeaderCreator
@@ -1677,12 +1780,6 @@ class SyclKernelIntHeaderCreator
     return true;
   }
 
-  // FIXME Remove this function when structs are replaced by their fields
-  bool handleStructType(FieldDecl *FD, QualType FieldTy) final {
-    addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout);
-    return true;
-  }
-
   bool handleScalarType(FieldDecl *FD, QualType FieldTy) final {
     addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout);
     return true;
@@ -1700,6 +1797,11 @@ class SyclKernelIntHeaderCreator
     return true;
   }
 
+  bool handleSyclHalfType(FieldDecl *FD, QualType FieldTy) final {
+    addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout);
+    return true;
+  }
+
   bool enterField(const CXXRecordDecl *RD, FieldDecl *FD) final {
     CurOffset += SemaRef.getASTContext().getFieldOffset(FD) / 8;
     return true;
@@ -1740,8 +1842,8 @@ class SyclKernelIntHeaderCreator
     CurOffset -= ArraySize;
     return true;
   }
-
   using SyclKernelFieldHandler::handleScalarType;
+  using SyclKernelFieldHandler::handleSyclHalfType;
   using SyclKernelFieldHandler::handleSyclSamplerType;
   using SyclKernelFieldHandler::leaveArray;
 };
@@ -1793,8 +1895,8 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc,
       StableName);
 
   ConstructingOpenCLKernel = true;
-  VisitRecordFields(KernelObj->fields(), checker, kernel_decl, kernel_body,
-                    int_header);
+  VisitRecordBases(KernelObj, checker, kernel_decl, kernel_body, int_header);
+  VisitRecordFields(KernelObj, checker, kernel_decl, kernel_body, int_header);
   ConstructingOpenCLKernel = false;
 }
 
@@ -2576,6 +2678,17 @@ bool Util::isSyclStreamType(const QualType &Ty) {
   return isSyclType(Ty, "stream");
 }
 
+bool Util::isSyclHalfType(const QualType &Ty) {
+  const StringRef &Name = "half";
+  std::array<DeclContextDesc, 5> Scopes = {
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "detail"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "half_impl"},
+      Util::DeclContextDesc{Decl::Kind::CXXRecord, Name}};
+  return matchQualifiedTypeName(Ty, Scopes);
+}
+
 bool Util::isSyclSpecConstantType(const QualType &Ty) {
   const StringRef &Name = "spec_constant";
   std::array<DeclContextDesc, 4> Scopes = {
diff --git a/clang/test/CodeGenSYCL/accessor_inheritance.cpp b/clang/test/CodeGenSYCL/accessor_inheritance.cpp
new file mode 100644
index 0000000000000..e197c339c1251
--- /dev/null
+++ b/clang/test/CodeGenSYCL/accessor_inheritance.cpp
@@ -0,0 +1,94 @@
+// RUN:  %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+#include <sycl.hpp>
+
+struct Base {
+  int A, B;
+  cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> AccField;
+};
+
+struct Captured : Base,
+                  cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> {
+  int C;
+};
+
+int main() {
+  Captured Obj;
+  cl::sycl::kernel_single_task<class kernel>(
+      [=]() {
+        Obj.use();
+      });
+  return 0;
+}
+
+// Check kernel parameters
+// CHECK: %[[RANGE_TYPE:"struct.*cl::sycl::range"]]
+// CHECK: %[[ID_TYPE:"struct.*cl::sycl::id"]]
+// CHECK: define spir_kernel void @_ZTSZ4mainE6kernel
+// CHECK-SAME: i32 [[ARG_A:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: i32 [[ARG_B:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: i8 addrspace(1)* [[ACC1_DATA:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) align 4 [[ACC1_RANGE1:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) align 4 [[ACC1_RANGE2:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: %[[ID_TYPE]]* byval(%[[ID_TYPE]]) align 4 [[ACC1_ID:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: i8 addrspace(1)* [[ACC2_DATA:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) align 4 [[ACC2_RANGE1:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) align 4 [[ACC2_RANGE2:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: %[[ID_TYPE]]* byval(%[[ID_TYPE]]) align 4 [[ACC2_ID:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: i32 [[ARG_C:%[a-zA-Z0-9_]+]])
+
+// Allocas for kernel parameters
+// CHECK: [[ARG_A]].addr = alloca i32
+// CHECK: [[ARG_B]].addr = alloca i32
+// CHECK: [[ACC1_DATA]].addr = alloca i8 addrspace(1)*
+// CHECK: [[ACC2_DATA]].addr = alloca i8 addrspace(1)*
+// CHECK: [[ARG_C]].addr = alloca i32
+//
+// Lambda object alloca
+// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_]+]] = alloca %"class.{{.*}}.anon"
+//
+// Kernel argument stores
+// CHECK: store i32 [[ARG_A]], i32* [[ARG_A]].addr
+// CHECK: store i32 [[ARG_B]], i32* [[ARG_B]].addr
+// CHECK: store i8 addrspace(1)* [[ACC1_DATA]], i8 addrspace(1)** [[ACC1_DATA]].addr
+// CHECK: store i8 addrspace(1)* [[ACC2_DATA]], i8 addrspace(1)** [[ACC2_DATA]].addr
+// CHECK: store i32 [[ARG_C]], i32* [[ARG_C]].addr
+//
+// Check A and B scalar fields initialization
+// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class._ZTSZ4mainE3$_0.anon", %"class._ZTSZ4mainE3$_0.anon"* [[KERNEL_OBJ]], i32 0, i32 0
+// CHECK: [[BITCAST:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured* [[GEP]] to %struct{{.*}}Base*
+// CHECK: [[FIELD_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST]], i32 0, i32 0
+// CHECK: [[ARG_A_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32* [[ARG_A]].addr
+// CHECK: store i32 [[ARG_A_LOAD]], i32* [[FIELD_A]]
+// CHECK: [[FIELD_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST]], i32 0, i32 1
+// CHECK: [[ARG_B_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32* [[ARG_B]].addr
+// CHECK: store i32 [[ARG_B_LOAD]], i32* [[FIELD_B]]
+//
+// Check accessors initialization
+// CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST]], i32 0, i32 2
+// CHECK: [[ACC1_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[ACC_FIELD]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
+// Default constructor call
+// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0EEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST]])
+// CHECK: [[BITCAST1:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured* [[GEP]] to i8*
+// CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8* [[BITCAST1]], i64 20
+// CHECK: [[BITCAST2:%[a-zA-Z0-9_]+]] = bitcast i8* [[GEP1]] to %"class{{.*}}cl::sycl::accessor"*
+// CHECK: [[ACC2_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[BITCAST2]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
+// Default constructor call
+// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0EEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC2_AS_CAST]])
+
+// CHECK C field initialization
+// CHECK: [[FIELD_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Captured, %struct{{.*}}Captured* [[GEP]], i32 0, i32 2
+// CHECK: [[ARG_C_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32* [[ARG_C]].addr
+// CHECK: store i32 [[ARG_C_LOAD]], i32* [[FIELD_C]]
+//
+// Check __init method calls
+// CHECK: [[GEP2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class._ZTSZ4mainE3$_0.anon", %"class._ZTSZ4mainE3$_0.anon"* [[KERNEL_OBJ]], i32 0, i32 0
+// CHECK: [[BITCAST3:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured* [[GEP2]] to %struct{{.*}}Base*
+// CHECK: [[ACC1_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST3]], i32 0, i32 2
+// CHECK: [[ACC1_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)** [[ACC1_DATA]].addr
+// CHECK: [[ACC1_AS_CAST1:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[ACC1_FIELD]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
+// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST1]], i8 addrspace(1)* [[ACC1_DATA_LOAD]]
+//
+// CHECK: [[GEP3:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class._ZTSZ4mainE3$_0.anon", %"class._ZTSZ4mainE3$_0.anon"* [[KERNEL_OBJ]], i32 0, i32 0
+// CHECK: [[ACC2_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)** [[ACC2_DATA]].addr
+// CHECK: [[AS_CAST_CAPTURED:%[a-zA-Z0-9_]+]] = addrspacecast %struct{{.*}}Captured* [[GEP3]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
+// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[AS_CAST_CAPTURED]], i8 addrspace(1)* [[ACC2_DATA_LOAD]]
diff --git a/clang/test/CodeGenSYCL/inheritance.cpp b/clang/test/CodeGenSYCL/inheritance.cpp
new file mode 100644
index 0000000000000..4ac785336fb39
--- /dev/null
+++ b/clang/test/CodeGenSYCL/inheritance.cpp
@@ -0,0 +1,84 @@
+// RUN:  %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+
+#include <sycl.hpp>
+
+class second_base {
+public:
+  int e;
+};
+
+class InnerFieldBase {
+public:
+  int d;
+};
+class InnerField : public InnerFieldBase {
+  int c;
+};
+
+struct base {
+public:
+  int b;
+  InnerField obj;
+};
+
+struct derived : base, second_base {
+  int a;
+
+  void operator()() {
+  }
+};
+
+int main() {
+  cl::sycl::queue q;
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    derived f{};
+    cgh.single_task(f);
+  });
+
+  return 0;
+}
+
+// Check kernel paramters
+// CHECK: define spir_kernel void @{{.*}}derived(i32 %_arg_b, i32 %_arg_d, i32 %_arg_c, i32 %_arg_e, i32 %_arg_a)
+
+// Check alloca for kernel paramters
+// CHECK: %[[ARG_B:[a-zA-Z0-9_.]+]] = alloca i32, align 4
+// CHECK: %[[ARG_D:[a-zA-Z0-9_.]+]] = alloca i32, align 4
+// CHECK: %[[ARG_C:[a-zA-Z0-9_.]+]] = alloca i32, align 4
+// CHECK: %[[ARG_E:[a-zA-Z0-9_.]+]] = alloca i32, align 4
+// CHECK: %[[ARG_A:[a-zA-Z0-9_.]+]] = alloca i32, align 4
+
+// Check alloca for local functor object
+// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = alloca %struct.{{.*}}.derived, align 4
+
+// Initialize field 'b'
+// CHECK: %[[BITCAST1:[0-9]+]] = bitcast %struct.{{.*}}.derived* %[[LOCAL_OBJECT]] to %struct.{{.*}}.base*
+// CHECK: %[[GEP_B:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.base, %struct.{{.*}}.base* %[[BITCAST1]], i32 0, i32 0
+// CHECK: %[[LOAD_B:[0-9]+]] = load i32, i32* %[[ARG_B]], align 4
+// CHECK: store i32 %[[LOAD_B]], i32* %[[GEP_B]], align 4
+
+// Initialize field 'd'
+// CHECK: %[[GEP_OBJ:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.base, %struct.{{.*}}.base* %[[BITCAST1]], i32 0, i32 1
+// CHECK: %[[BITCAST2:[0-9]+]] = bitcast %class.{{.*}}.InnerField* %[[GEP_OBJ]] to %class.{{.*}}.InnerFieldBase*
+// CHECK: %[[GEP_D:[a-zA-Z0-9]+]] = getelementptr inbounds %class.{{.*}}.InnerFieldBase, %class.{{.*}}.InnerFieldBase* %[[BITCAST2]], i32 0, i32 0
+// CHECK: %[[LOAD_D:[0-9]+]] = load i32, i32* %[[ARG_D]], align 4
+// CHECK: store i32 %[[LOAD_D]], i32* %[[GEP_D]], align 4
+
+// Initialize field 'c'
+// CHECK: %[[GEP_C:[a-zA-Z0-9]+]] = getelementptr inbounds %class.{{.*}}.InnerField, %class.{{.*}}.InnerField* %[[GEP_OBJ]], i32 0, i32 1
+// CHECK: %[[LOAD_C:[0-9]+]] = load i32, i32* %[[ARG_C]], align 4
+// CHECK: store i32 %[[LOAD_C]], i32* %[[GEP_C]], align 4
+
+// Initialize field 'e'
+// CHECK: %[[BITCAST3:[0-9]+]] = bitcast %struct.{{.*}}.derived* %[[LOCAL_OBJECT]] to i8*
+// CHECK: %[[GEP_DERIVED:[a-zA-Z0-9]+]] = getelementptr inbounds i8, i8* %[[BITCAST3]], i64 12
+// CHECK: %[[BITCAST4:[0-9]+]] = bitcast i8* %[[GEP_DERIVED]] to %class.{{.*}}.second_base*
+// CHECK: %[[GEP_E:[a-zA-Z0-9]+]] = getelementptr inbounds %class.{{.*}}.second_base, %class.{{.*}}.second_base* %[[BITCAST4]], i32 0, i32 0
+// CHECK: %[[LOAD_E:[0-9]+]] = load i32, i32* %[[ARG_E]], align 4
+// CHECK: store i32 %[[LOAD_E]], i32* %[[GEP_E]], align 4
+
+// Initialize field 'a'
+// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.derived, %struct.{{.*}}.derived* %[[LOCAL_OBJECT]], i32 0, i32 2
+// CHECK: %[[LOAD_A:[0-9]+]] = load i32, i32* %[[ARG_A]], align 4
+// CHECK: store i32 %[[LOAD_A]], i32* %[[GEP_A]], align 4
diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp
index 58d0c3addcd8c..1c766b2dccd3f 100644
--- a/clang/test/CodeGenSYCL/integration_header.cpp
+++ b/clang/test/CodeGenSYCL/integration_header.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -emit-llvm
 // RUN: FileCheck -input-file=%t.h %s
 //
 // CHECK: #include <CL/sycl/detail/kernel_desc.hpp>
@@ -28,9 +28,11 @@
 // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
 // CHECK-NEXT:   //--- _ZTSZ4mainE12first_kernel
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 0 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 4 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 6112, 16 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_sampler, 8, 32 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 1, 4 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 8 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 12 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 6112, 24 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_sampler, 8, 40 },
 // CHECK-EMPTY:
 // CHECK-NEXT:   //--- _ZTSN16second_namespace13second_kernelIcEE
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 0 },
@@ -46,12 +48,15 @@
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 0 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 6112, 4 },
 // CHECK-EMPTY:
-// CHECK-NEXT:   //--- _ZTSZ4mainE16accessor_in_base
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 64, 0 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 8 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 24 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 40 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 52 },
+// CHECK-NEXT:  //--- _ZTSZ4mainE16accessor_in_base
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 0 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 4 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_accessor, 4062, 8 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 20 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_accessor, 4062, 24 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 36 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_accessor, 4062, 40 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_accessor, 4062, 52 },
 // CHECK-EMPTY:
 // CHECK-NEXT: };
 //
@@ -116,15 +121,13 @@ int main() {
       acc2;
   int i = 13;
   cl::sycl::sampler smplr;
-  // TODO: Uncomemnt when structures in kernel arguments are correctly processed
-  //       by SYCL compiler
-  /*  struct {
+  struct {
     char c;
     int i;
   } test_s;
-  test_s.c = 14;*/
+  test_s.c = 14;
   kernel_single_task<class first_kernel>([=]() {
-    if (i == 13 /*&& test_s.c == 14*/) {
+    if (i == 13 && test_s.c == 14) {
 
       acc1.use();
       acc2.use();
@@ -151,10 +154,9 @@ int main() {
       }
   });
 
-  // FIXME: We cannot use the member-capture because all the handlers except the
-  // integration header handler in SemaSYCL don't handle base types right.
   accessor_in_base::captured c;
-  kernel_single_task<class accessor_in_base>([c]() {
+  kernel_single_task<class accessor_in_base>([=]() {
+    c.use();
   });
 
   return 0;
diff --git a/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp b/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp
old mode 100755
new mode 100644
diff --git a/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp b/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp
index 141191219b4dc..f5f679f7d3650 100644
--- a/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp
@@ -1,7 +1,5 @@
 // RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only
 // RUN: FileCheck -input-file=%t.h %s
-// XFAIL for now due to : https://github.com/intel/llvm/issues/2018
-// XFAIL: *
 
 // This test checks the integration header when kernel argument
 // is a struct containing an Accessor array.
@@ -22,7 +20,6 @@
 // CHECK: static constexpr
 // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
 // CHECK-NEXT: //--- _ZTSZ4mainE8kernel_C
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 24, 0 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 0 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 12 },
 // CHECK-EMPTY:
diff --git a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp
index ae476edf08c2e..1b1b25dcd3ff4 100644
--- a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp
@@ -1,5 +1,4 @@
 // RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -fsycl-int-header=%t.h -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
-// XFAIL: *
 
 // This test checks a kernel with struct parameter that contains an Accessor array.
 
@@ -29,7 +28,6 @@ int main() {
 
 // CHECK kernel_C parameters
 // CHECK: define spir_kernel void @{{.*}}kernel_C
-// CHECK-SAME: %struct.{{.*}}.struct_acc_t* byval(%struct.{{.*}}.struct_acc_t) align 4 [[STRUCT:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: i32 addrspace(1)* [[MEM_ARG1:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
 // CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
@@ -54,32 +52,26 @@ int main() {
 // CHECK: [[MEM_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
 // CHECK: [[OFFSET2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id"
 
-// Check init of local struct
-// CHECK: [[L_STRUCT_ADDR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
-// CHECK: [[MEMCPY_DST:%[0-9a-zA-Z_]+]] = bitcast %struct.{{.*}}struct_acc_t* [[L_STRUCT_ADDR]] to i8*
-// CHECK: [[MEMCPY_SRC:%[0-9a-zA-Z_]+]] = bitcast %struct.{{.*}}struct_acc_t* %{{[0-9a-zA-Z_]+}} to i8*
-// CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 [[MEMCPY_DST]], i8* align 4 [[MEMCPY_SRC]], i64 24, i1 false)
-
-// Check accessor array GEP for member_acc[0]
-// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
-// CHECK: [[MEMBER1:%[a-zA-Z_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[ACCESSOR_ARRAY1]], i32 0, i32 0
-// CHECK: [[Z0:%[a-zA-Z0-9_]*]] = getelementptr inbounds [2 x %"class.{{.*}}.cl::sycl::accessor"], [2 x %"class.{{.*}}.cl::sycl::accessor"]* [[MEMBER1]], i64 0, i64 0
-
-// Check load from kernel pointer argument alloca
-// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr{{[0-9]*}}
+// Check loop which calls the default constructor for each element of accessor array is emitted.
+// CHECK: [[GEP_LAMBDA:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: [[GEP_MEMBER_ACC:%[a-zA-Z_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[GEP_LAMBDA]], i32 0, i32 0
+// CHECK: [[ARRAY_BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], [2 x [[ACCESSOR]]]* [[GEP_MEMBER_ACC]], i64 0, i64 0
+// CHECK: [[ARRAY_END:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* [[ARRAY_BEGIN]], i64 2
+// CHECK: br label %arrayctor.loop
+// CHECK: arrayctor.loop:
 
 // Check acc[0] __init method call
-// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[Z0]] to %"class{{.*}}accessor" addrspace(4)*
-// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACC_CAST1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]])
-
-// Check accessor array GEP for member_acc[1]
-// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
-// CHECK: [[MEMBER2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[ACCESSOR_ARRAY2]], i32 0, i32 0
-// CHECK: [[Z1:%[a-zA-Z0-9_]*]] = getelementptr inbounds [2 x %"class.{{.*}}.cl::sycl::accessor"], [2 x %"class.{{.*}}.cl::sycl::accessor"]* [[MEMBER2]], i64 0, i64 1
-
-// Check load from kernel pointer argument alloca
-// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr{{[0-9]*}}
+// CHECK: [[GEP_LAMBDA1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: [[GEP_MEMBER_ACC1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[GEP_LAMBDA1]], i32 0, i32 0
+// CHECK: [[ARRAY_IDX1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[GEP_MEMBER_ACC1]], i64 0, i64 0
+// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr
+// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast [[ACCESSOR]]* [[ARRAY_IDX1]] to [[ACCESSOR]] addrspace(4)*
+// CHECK: call spir_func void @{{.*}}__init{{.*}}([[ACCESSOR]] addrspace(4)* [[ACC_CAST1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]])
 
 // Check acc[1] __init method call
-// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[Z1]] to %"class{{.*}}accessor" addrspace(4)*
-// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACC_CAST2]], i32 addrspace(1)* [[MEM_LOAD2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]])
+// CHECK: [[GEP_LAMBDA2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: [[GEP_MEMBER_ACC2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[GEP_LAMBDA2]], i32 0, i32 0
+// CHECK: [[ARRAY_IDX2:%[a-zA-Z0-9_]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[GEP_MEMBER_ACC2]], i64 0, i64 1
+// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr
+// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast [[ACCESSOR]]* [[ARRAY_IDX2]] to [[ACCESSOR]] addrspace(4)*
+// CHECK: call spir_func void @{{.*}}__init{{.*}}([[ACCESSOR]] addrspace(4)* [[ACC_CAST2]], i32 addrspace(1)* [[MEM_LOAD2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]])
diff --git a/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp b/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp
index 49fd34d3206e5..d4a5c8d5995a0 100755
--- a/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp
@@ -1,6 +1,5 @@
 // RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only
 // RUN: FileCheck -input-file=%t.h %s
-
 // This test checks the integration header generated for a kernel
 // with an argument that is a POD array.
 
diff --git a/clang/test/CodeGenSYCL/sampler.cpp b/clang/test/CodeGenSYCL/sampler.cpp
index 749b5a0bfdaa9..39755aa6b9aec 100644
--- a/clang/test/CodeGenSYCL/sampler.cpp
+++ b/clang/test/CodeGenSYCL/sampler.cpp
@@ -12,11 +12,28 @@
 // CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler" addrspace(4)* [[GEPCAST]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]])
 //
 
-// CHECK: define spir_kernel void @{{[a-zA-Z0-9_]+}}(%struct{{.*}}sampler_wrapper{{.*}} %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG_WRAPPED:%[a-zA-Z0-9_]+]])
+// CHECK: define spir_kernel void @{{[a-zA-Z0-9_]+}}(%opencl.sampler_t addrspace(2)* [[SAMPLER_ARG_WRAPPED:%[a-zA-Z0-9_]+]], i32 [[ARG_A:%[a-zA-Z0-9_]+]])
+
+// Check alloca
 // CHECK: [[SAMPLER_ARG_WRAPPED]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8
+// CHECK: [[ARG_A]].addr = alloca i32, align 4
+// CHECK: [[LAMBDA:%[0-9]+]] = alloca %"class.{{.*}}.anon.0", align 8
+
+// Check argument store
 // CHECK: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG_WRAPPED]], %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG_WRAPPED]].addr, align 8
-// CHECK: [[LOAD_SAMPLER_ARG_WRAPPED:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG_WRAPPED]].addr, align 8
-// CHECK: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler" addrspace(4)* {{.*}}, %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG_WRAPPED]])
+// CHECK: store i32 [[ARG_A]], i32* [[ARG_A]].addr, align 4
+
+// Initialize 'a'
+// CHECK: [[GEP_LAMBDA:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon.0", %"class.{{.*}}.anon.0"* [[LAMBDA]], i32 0, i32 0
+// CHECK: [[GEP_A:%[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.sampler_wrapper, %struct.{{.*}}.sampler_wrapper* [[GEP_LAMBDA]], i32 0, i32 1
+// CHECK: [[LOAD_A:%[0-9]+]] = load i32, i32* [[ARG_A]].addr, align 4
+// CHECK: store i32 [[LOAD_A]], i32* [[GEP_A]], align 8
+
+// Initialize wrapped sampler 'smpl'
+// CHECK: [[GEP_LAMBDA_0:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon.0", %"class.{{.*}}.anon.0"* %0, i32 0, i32 0
+// CHECK: [[GEP_SMPL:%[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.sampler_wrapper, %struct.{{.*}}.sampler_wrapper* [[GEP_LAMBDA_0]], i32 0, i32 0
+// CHECK: [[LOAD_SMPL:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG_WRAPPED]].addr, align 8
+// CHECK: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler" addrspace(4)* {{.*}}, %opencl.sampler_t addrspace(2)* [[LOAD_SMPL]])
 //
 #include "sycl.hpp"
 
diff --git a/clang/test/CodeGenSYCL/struct_kernel_param.cpp b/clang/test/CodeGenSYCL/struct_kernel_param.cpp
index 8e6fbcec309dd..896cf8c54ec84 100644
--- a/clang/test/CodeGenSYCL/struct_kernel_param.cpp
+++ b/clang/test/CodeGenSYCL/struct_kernel_param.cpp
@@ -4,11 +4,15 @@
 // CHECK:     const kernel_param_desc_t kernel_signatures[] = {
 // CHECK-NEXT:  //--- _ZTSZZ5test0vENK3$_0clERN2cl4sycl7handlerEE8MyKernel
 // CHECK-NEXT:  { kernel_param_kind_t::kind_accessor, 4062, 0 },
-// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 24, 12 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 12 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 16 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 20 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 24 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 28 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 32 },
 // CHECK-EMPTY:
 // CHECK-NEXT:};
 
-
 // This test checks if compiler accepts structures as kernel parameters.
 
 #include "sycl.hpp"
diff --git a/clang/test/CodeGenSYCL/wrapped-accessor.cpp b/clang/test/CodeGenSYCL/wrapped-accessor.cpp
index 0cd651efc58f5..1380aebba234e 100644
--- a/clang/test/CodeGenSYCL/wrapped-accessor.cpp
+++ b/clang/test/CodeGenSYCL/wrapped-accessor.cpp
@@ -17,7 +17,6 @@
 // CHECK: static constexpr
 // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
 // CHECK-NEXT: //--- _ZTSZ4mainE14wrapped_access
-// CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 12, 0 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 0 },
 // CHECK-EMPTY:
 // CHECK-NEXT: };
diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp
index 5bd37447ce814..9e3efc6321096 100644
--- a/clang/test/SemaSYCL/Inputs/sycl.hpp
+++ b/clang/test/SemaSYCL/Inputs/sycl.hpp
@@ -37,6 +37,20 @@ enum class address_space : int {
 };
 } // namespace access
 
+namespace detail {
+namespace half_impl {
+struct half {
+#ifdef __SYCL_DEVICE_ONLY
+  _Float16 data;
+#else
+  char data[2];
+#endif
+};
+} // namespace half_impl
+} // namespace detail
+
+using half = detail::half_impl::half;
+
 template <int dim>
 struct range {
 };
diff --git a/clang/test/SemaSYCL/accessor_inheritance.cpp b/clang/test/SemaSYCL/accessor_inheritance.cpp
new file mode 100644
index 0000000000000..17dafe7b4acdd
--- /dev/null
+++ b/clang/test/SemaSYCL/accessor_inheritance.cpp
@@ -0,0 +1,65 @@
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s
+#include <sycl.hpp>
+
+struct Base {
+  int A, B;
+  cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> AccField;
+};
+
+struct Captured : Base,
+                  cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> {
+  int C;
+};
+
+int main() {
+  Captured Obj;
+  cl::sycl::kernel_single_task<class kernel>(
+      [=]() {
+        Obj.use();
+      });
+}
+
+// Check kernel parameters
+// CHECK: FunctionDecl {{.*}}kernel{{.*}} 'void (int, int, __global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int)'
+// CHECK: ParmVarDecl{{.*}} used _arg_A 'int'
+// CHECK: ParmVarDecl{{.*}} used _arg_B 'int'
+// CHECK: ParmVarDecl{{.*}} used _arg_AccField '__global char *'
+// CHECK: ParmVarDecl{{.*}} used _arg_AccField 'cl::sycl::range<1>'
+// CHECK: ParmVarDecl{{.*}} used _arg_AccField 'cl::sycl::range<1>'
+// CHECK: ParmVarDecl{{.*}} used _arg_AccField 'cl::sycl::id<1>'
+// CHECK: ParmVarDecl{{.*}} used _arg__base '__global char *'
+// CHECK: ParmVarDecl{{.*}} used _arg__base 'cl::sycl::range<1>'
+// CHECK: ParmVarDecl{{.*}} used _arg__base 'cl::sycl::range<1>'
+// CHECK: ParmVarDecl{{.*}} used _arg__base 'cl::sycl::id<1>'
+// CHECK: ParmVarDecl{{.*}} used _arg_C 'int'
+
+// Check lambda initialization
+// CHECK: VarDecl {{.*}} used '(lambda at {{.*}}accessor_inheritance.cpp
+// CHECK-NEXT: InitListExpr {{.*}}
+// CHECK-NEXT: InitListExpr {{.*}} 'Captured'
+// CHECK-NEXT: InitListExpr {{.*}} 'Base'
+// 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 {{.*}} 'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read>':'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>' 'void () noexcept'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read>':'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>' 'void () noexcept'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_C' 'int'
+
+// Check __init calls
+// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK-NEXT: MemberExpr {{.*}} .__init
+// CHECK-NEXT: MemberExpr {{.*}} .AccField
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'Base' lvalue <DerivedToBase (Base)>
+// CHECK-NEXT: MemberExpr {{.*}} 'Captured' 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 *'
+
+// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK-NEXT: MemberExpr{{.*}} lvalue .__init
+// CHECK-NEXT: MemberExpr{{.*}}'Captured' 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 *'
diff --git a/clang/test/SemaSYCL/array-kernel-param.cpp b/clang/test/SemaSYCL/array-kernel-param.cpp
old mode 100755
new mode 100644
index c8bdb390467a1..42c1ef42c2256
--- a/clang/test/SemaSYCL/array-kernel-param.cpp
+++ b/clang/test/SemaSYCL/array-kernel-param.cpp
@@ -20,7 +20,7 @@ int main() {
   Accessor acc[2];
   int a[2];
   struct struct_acc_t {
-    Accessor member_acc[4];
+    Accessor member_acc[2];
   } struct_acc;
 
   a_kernel<class kernel_A>(
@@ -69,33 +69,27 @@ int main() {
 // CHECK: ImplicitCastExpr
 // CHECK: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
 
-// Correct and enable after struct members are extracted into separate parameters
-// C HECK kernel_C parameters
-// C HECK: FunctionDecl {{.*}}kernel_C{{.*}} 'void (struct {{.*}}, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
-// C HECK-NEXT: ParmVarDecl {{.*}} 'struct {{.*}}'
-// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
-// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
-// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
-// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
-// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
-// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
-// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
-// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
-// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
-// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
-// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
-// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
-// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
-// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
-// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
-// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
+// Check kernel_C parameters
+// CHECK: FunctionDecl {{.*}}kernel_C{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}array-kernel-param.cpp:37:7)' cinit
+// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp:37:7)'
+// CHECK-NEXT: InitListExpr {{.*}} 'struct_acc_t'
+// CHECK-NEXT: InitListExpr {{.*}} 'Accessor [2]'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'Accessor [2]'
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'Accessor [2]'
 
-// C HECK that four accessor init functions are called
-// C HECK: CXXMemberCallExpr {{.*}} 'void'
-// C HECK-NEXT: MemberExpr {{.*}}__init
-// C HECK: CXXMemberCallExpr {{.*}} 'void'
-// C HECK-NEXT: MemberExpr {{.*}}__init
-// C HECK: CXXMemberCallExpr {{.*}} 'void'
-// C HECK-NEXT: MemberExpr {{.*}}__init
-// C HECK: CXXMemberCallExpr {{.*}} 'void'
-// C HECK-NEXT: MemberExpr {{.*}}__init
+// Check __init functions are called
+// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK-NEXT: MemberExpr {{.*}}__init
+// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK-NEXT: MemberExpr {{.*}}__init
diff --git a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp
index 82cd21bf01552..2d04e0f453579 100644
--- a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp
+++ b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp
@@ -65,15 +65,15 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
 
 // Check kernel parameters
-// CHECK: {{.*}}kernel_struct{{.*}} 'void (test_struct)'
-// CHECK: ParmVarDecl {{.*}} used _arg_ 'test_struct'
+// CHECK: {{.*}}kernel_struct{{.*}} 'void (int)'
+// CHECK: ParmVarDecl {{.*}} used _arg_data 'int'
 
 // Check that lambda field of struct type is initialized
 // CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})'
-// CHECK-NEXT: InitListExpr
-// CHECK-NEXT: CXXConstructExpr {{.*}}'test_struct'{{.*}}void (const test_struct &)
-// CHECK-NEXT: ImplicitCastExpr {{.*}}'const test_struct' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'test_struct' lvalue ParmVar {{.*}} '_arg_' 'test_struct'
+// CHECK-NEXT: InitListExpr {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})'
+// CHECK-NEXT: InitListExpr {{.*}}'test_struct'{{.*}}
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_data' 'int'
 
 // Check kernel parameters
 // CHECK: {{.*}}kernel_pointer{{.*}} 'void (__global int *, __global int *)'
diff --git a/clang/test/SemaSYCL/fake-accessors.cpp b/clang/test/SemaSYCL/fake-accessors.cpp
index 24d36a6ba54b6..a1357083dd2ce 100644
--- a/clang/test/SemaSYCL/fake-accessors.cpp
+++ b/clang/test/SemaSYCL/fake-accessors.cpp
@@ -51,6 +51,6 @@ int main() {
         });
   return 0;
 }
-// CHECK: fake_accessors{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
-// CHECK: accessor_typedef{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
-// CHECK: accessor_alias{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
+// CHECK: fake_accessors{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int, int)
+// CHECK: accessor_typedef{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int, int)
+// CHECK: accessor_alias{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int, int)
diff --git a/clang/test/SemaSYCL/half-kernel-arg.cpp b/clang/test/SemaSYCL/half-kernel-arg.cpp
new file mode 100644
index 0000000000000..13f89f38a3d96
--- /dev/null
+++ b/clang/test/SemaSYCL/half-kernel-arg.cpp
@@ -0,0 +1,23 @@
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s
+
+// This test checks that compiler generates correct initialization for arguments
+// that have cl::sycl::half type inside the OpenCL kernel
+
+#include <sycl.hpp>
+
+int main() {
+  cl::sycl::half HostHalf;
+  cl::sycl::kernel_single_task<class kernel_half>(
+      [=]() {
+        cl::sycl::half KernelHalf = HostHalf;
+      });
+}
+
+// CHECK: {{.*}}kernel_half{{.*}} 'void (cl::sycl::half)'
+// CHECK: ParmVarDecl {{.*}} used _arg_ 'cl::sycl::half':'cl::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 {{.*}}'cl::sycl::detail::half_impl::half'{{.*}}
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::detail::half_impl::half'
+// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::half':'cl::sycl::detail::half_impl::half' lvalue ParmVar {{.*}} '_arg_' 'cl::sycl::half':'cl::sycl::detail::half_impl::half'
diff --git a/clang/test/SemaSYCL/inheritance.cpp b/clang/test/SemaSYCL/inheritance.cpp
new file mode 100644
index 0000000000000..ff0b263449a35
--- /dev/null
+++ b/clang/test/SemaSYCL/inheritance.cpp
@@ -0,0 +1,70 @@
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s
+
+#include <sycl.hpp>
+
+class second_base {
+public:
+  int e;
+};
+
+class InnerFieldBase {
+public:
+  int d;
+};
+class InnerField : public InnerFieldBase {
+  int c;
+};
+
+struct base {
+public:
+  int b;
+  InnerField obj;
+};
+
+struct derived : base, second_base {
+  int a;
+
+  void operator()() {
+  }
+};
+
+int main() {
+  cl::sycl::queue q;
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    derived f{};
+    cgh.single_task(f);
+  });
+
+  return 0;
+}
+
+// Check declaration of the kernel
+// CHECK: derived{{.*}} 'void (int, int, int, int, int)'
+
+// Check parameters of the kernel
+// CHECK: ParmVarDecl {{.*}} used _arg_b 'int'
+// CHECK: ParmVarDecl {{.*}} used _arg_d 'int'
+// CHECK: ParmVarDecl {{.*}} used _arg_c 'int'
+// CHECK: ParmVarDecl {{.*}} used _arg_e 'int'
+// CHECK: ParmVarDecl {{.*}} used _arg_a 'int'
+
+// Check initializers for derived and base classes.
+// Each class has it's own initializer list
+// Base classes should be initialized first.
+// CHECK: VarDecl {{.*}} derived 'derived' cinit
+// CHECK-NEXT: InitListExpr {{.*}} 'derived'
+// CHECK-NEXT: InitListExpr {{.*}} 'base'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_b' 'int'
+// CHECK-NEXT: InitListExpr {{.*}} 'InnerField'
+// CHECK-NEXT: InitListExpr {{.*}} 'InnerFieldBase'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_d' 'int'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_c' 'int'
+// CHECK-NEXT: InitListExpr {{.*}} 'second_base'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_e' 'int'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_a' 'int'
diff --git a/clang/test/SemaSYCL/wrapped-accessor.cpp b/clang/test/SemaSYCL/wrapped-accessor.cpp
index 83bb3ff2448fb..1052b4ac24e0f 100644
--- a/clang/test/SemaSYCL/wrapped-accessor.cpp
+++ b/clang/test/SemaSYCL/wrapped-accessor.cpp
@@ -23,10 +23,9 @@ int main() {
 }
 
 // Check declaration of the kernel
-// CHECK: wrapped_access{{.*}} 'void (AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
+// CHECK: wrapped_access{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
 
 // Check parameters of the kernel
-// CHECK: ParmVarDecl {{.*}} used _arg_ 'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>':'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>'
 // CHECK: ParmVarDecl {{.*}} used _arg_accessor '__global int *'
 // CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>'
 // CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>'
@@ -35,10 +34,9 @@ int main() {
 // Check that wrapper object itself is initialized with corresponding kernel
 // argument
 // CHECK: VarDecl {{.*}}'(lambda at {{.*}}wrapped-accessor.cpp{{.*}})'
-// CHECK-NEXT: InitListExpr
-// CHECK-NEXT: CXXConstructExpr {{.*}}AccWrapper<cl::sycl::accessor
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>':'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>' lvalue ParmVar {{.*}} '_arg_' 'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>':'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>'
+// CHECK-NEXT: InitListExpr {{.*}}'(lambda at {{.*}}wrapped-accessor.cpp{{.*}})'
+// CHECK-NEXT: InitListExpr {{.*}}'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>'
+// CHECK-NEXT: CXXConstructExpr {{.*}}'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>':'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>' 'void () noexcept'
 
 // Check that accessor field of the wrapper object is initialized using __init method
 // CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
diff --git a/sycl/test/array_param/array-kernel-param-nested-run.cpp b/sycl/test/array_param/array-kernel-param-nested-run.cpp
index 28b9469cda89a..b8fcd8c9dfc1f 100755
--- a/sycl/test/array_param/array-kernel-param-nested-run.cpp
+++ b/sycl/test/array_param/array-kernel-param-nested-run.cpp
@@ -5,7 +5,6 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
-// XFAIL: *
 
 #include <CL/sycl.hpp>
 #include <iostream>
diff --git a/sycl/test/basic_tests/accessor/accessor.cpp b/sycl/test/basic_tests/accessor/accessor.cpp
index 84286c6117eff..ebb99a0f7abbe 100644
--- a/sycl/test/basic_tests/accessor/accessor.cpp
+++ b/sycl/test/basic_tests/accessor/accessor.cpp
@@ -30,6 +30,16 @@ struct IdxID3 {
   operator sycl::id<3>() { return sycl::id<3>(x, y, z); }
 };
 
+template <typename T>
+using AccAlias =
+    cl::sycl::accessor<T, 1, cl::sycl::access::mode::write,
+                       cl::sycl::access::target::global_buffer>;
+template <typename T>
+struct InheritedAccessor : public AccAlias<T> {
+
+  using AccAlias<T>::AccAlias;
+};
+
 template <typename Acc> struct AccWrapper { Acc accessor; };
 
 template <typename Acc1, typename Acc2> struct AccsWrapper {
@@ -493,4 +503,41 @@ int main() {
       return 1;
     }
   }
+  {
+    try {
+      int data = -1;
+      int cnst = 399;
+
+      {
+        sycl::buffer<int, 1> A(&cnst, sycl::range<1>(1));
+        sycl::buffer<int, 1> B(&cnst, sycl::range<1>(1));
+        sycl::buffer<int, 1> C(&data, sycl::range<1>(1));
+
+        sycl::queue queue;
+        queue.submit([&](sycl::handler &cgh) {
+          sycl::accessor<int, 1, sycl::access::mode::write,
+                         sycl::access::target::global_buffer>
+              AccA(A, cgh);
+          sycl::accessor<int, 1, sycl::access::mode::read,
+                         sycl::access::target::constant_buffer>
+              AccB(B, cgh);
+          InheritedAccessor<int> AccC(C, cgh);
+          cgh.single_task<class acc_base>([=]() {
+            AccC[0] = AccA[0] + AccB[0];
+          });
+        });
+
+#ifndef simplification_test
+        auto host_acc = C.get_access<sycl::access::mode::read>();
+#else
+        sycl::host_accessor host_acc(C, sycl::read_only);
+#endif
+        assert(host_acc[0] == 798);
+      }
+
+    } catch (sycl::exception e) {
+      std::cout << "SYCL exception caught: " << e.what();
+      return 1;
+    }
+  }
 }
diff --git a/sycl/test/functor/functor_inheritance.cpp b/sycl/test/functor/functor_inheritance.cpp
new file mode 100644
index 0000000000000..3a3d5218ef6ec
--- /dev/null
+++ b/sycl/test/functor/functor_inheritance.cpp
@@ -0,0 +1,61 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -o %t.out %s
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+#include <CL/sycl.hpp>
+#include <cassert>
+
+constexpr auto sycl_read_write = cl::sycl::access::mode::read_write;
+constexpr auto sycl_global_buffer = cl::sycl::access::target::global_buffer;
+
+struct SecondBase {
+  SecondBase(int _E) : E(_E) {}
+  int E;
+};
+
+struct InnerFieldBase {
+  InnerFieldBase(int _D) : D(_D) {}
+  int D;
+};
+
+struct InnerField : public InnerFieldBase {
+  InnerField(int _C, int _D) : C(_C), InnerFieldBase(_D) {}
+  int C;
+};
+
+struct Base {
+  Base(int _B, int _C, int _D) : B(_B), InnerObj(_C, _D) {}
+  int B;
+  InnerField InnerObj;
+};
+
+struct Derived : public Base, public SecondBase {
+  Derived(
+      int _A, int _B, int _C, int _D, int _E,
+      cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> &_Acc)
+      : A(_A), Acc(_Acc), /*Out(_Out),*/ Base(_B, _C, _D), SecondBase(_E) {}
+  void operator()() {
+    Acc[0] = this->A + this->B + this->InnerObj.C + this->InnerObj.D + this->E;
+  }
+
+  int A;
+  cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> Acc;
+};
+
+int main() {
+  int A[] = {10};
+  {
+    cl::sycl::queue Q;
+    cl::sycl::buffer<int, 1> Buf(A, 1);
+
+    Q.submit([&](cl::sycl::handler &cgh) {
+      auto Acc = Buf.get_access<sycl_read_write, sycl_global_buffer>(cgh);
+      Derived F = {1, 2, 3, 4, 5, Acc /*, Out*/};
+      cgh.single_task(F);
+    });
+  }
+  assert(A[0] == 15);
+  return 0;
+}