diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index b94bc75b081b2..b540d77b84d40 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1232,6 +1232,13 @@ def SYCLKernel : InheritableAttr {
   let Documentation = [SYCLKernelDocs];
 }
 
+def SYCLSpecialClass: InheritableAttr {
+  let Spellings = [Clang<"sycl_special_class">];
+  let Subjects = SubjectList<[CXXRecord]>;
+  let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
+  let Documentation = [SYCLSpecialClassDocs];
+}
+
 // Marks functions which must not be vectorized via horizontal SIMT widening,
 // e.g. because the function is already vectorized. Used to mark SYCL
 // explicit SIMD kernels and functions.
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 09949d163c627..e37a3c7a72a57 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -405,6 +405,26 @@ The SYCL kernel in the previous code sample meets these expectations.
   }];
 }
 
+def SYCLSpecialClassDocs : Documentation {
+  let Category = DocCatStmt;
+  let Content = [{
+The ``__attribute__((sycl_special_class))`` attribute is used in SYCL
+headers to indicate that a class or a struct needs additional handling when
+it is passed from host to device. Please note that this is an attribute that is
+used for internal implementation and not intended to be used by external users.
+It is used for ``accessor``, ``sampler`` , or ``stream`` classes.
+The types that own this attribute are excluded from device-copyable and other
+type-legalization steps.
+
+.. code-block:: c++
+  class __attribute__((sycl_special_class)) accessor {
+
+private:
+  void __init() {}
+};
+  }];
+}
+
 def SYCLSimdDocs : Documentation {
   let Category = DocCatFunction;
   let Content = [{
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index ccdd618315575..1ff52eb786898 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -9782,6 +9782,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
   case ParsedAttr::AT_SYCLSimd:
     handleSimpleAttribute<SYCLSimdAttr>(S, D, AL);
     break;
+  case ParsedAttr::AT_SYCLSpecialClass:
+    handleSimpleAttribute<SYCLSpecialClassAttr>(S, D, AL);
+    break;
   case ParsedAttr::AT_SYCLDevice:
     handleSYCLDeviceAttr(S, D, AL);
     break;
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 8acf1962b4bc0..0b989e747a9cd 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -89,17 +89,7 @@ class Util {
     return DeclContextDesc{K, SR};
   }
 
-  /// Checks whether given clang type is a full specialization of the SYCL
-  /// accessor class.
-  static bool isSyclAccessorType(QualType Ty);
-
-  /// Checks whether given clang type is a full specialization of the SYCL
-  /// sampler class.
-  static bool isSyclSamplerType(QualType Ty);
-
-  /// Checks whether given clang type is a full specialization of the SYCL
-  /// stream class.
-  static bool isSyclStreamType(QualType Ty);
+  static bool isSyclSpecialType(QualType Ty);
 
   /// Checks whether given clang type is a full specialization of the SYCL
   /// half class.
@@ -1224,13 +1214,9 @@ class KernelObjVisitor {
     for (const auto &Base : Range) {
       QualType BaseTy = Base.getType();
       // Handle accessor class as base
-      if (Util::isSyclAccessorType(BaseTy)) {
-        (void)std::initializer_list<int>{
-            (Handlers.handleSyclAccessorType(Owner, Base, BaseTy), 0)...};
-      } else if (Util::isSyclStreamType(BaseTy)) {
-        // Handle stream class as base
+      if (Util::isSyclSpecialType(BaseTy)) {
         (void)std::initializer_list<int>{
-            (Handlers.handleSyclStreamType(Owner, Base, BaseTy), 0)...};
+            (Handlers.handleSyclSpecialType(Owner, Base, BaseTy), 0)...};
       } else
         // For all other bases, visit the record
         visitRecord(Owner, Base, BaseTy->getAsCXXRecordDecl(), BaseTy,
@@ -1311,16 +1297,12 @@ class KernelObjVisitor {
   template <typename... HandlerTys>
   void visitField(const CXXRecordDecl *Owner, FieldDecl *Field,
                   QualType FieldTy, HandlerTys &... Handlers) {
-    if (Util::isSyclAccessorType(FieldTy))
-      KF_FOR_EACH(handleSyclAccessorType, Field, FieldTy);
-    else if (Util::isSyclSamplerType(FieldTy))
-      KF_FOR_EACH(handleSyclSamplerType, Field, FieldTy);
+    if (Util::isSyclSpecialType(FieldTy))
+      KF_FOR_EACH(handleSyclSpecialType, 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))
-      KF_FOR_EACH(handleSyclStreamType, Field, FieldTy);
     else if (FieldTy->isStructureOrClassType()) {
       if (KF_FOR_EACH(handleStructType, Field, FieldTy)) {
         CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl();
@@ -1375,25 +1357,18 @@ class SyclKernelFieldHandlerBase {
   // Mark these virtual so that we can use override in the implementer classes,
   // despite virtual dispatch never being used.
 
-  // Accessor can be a base class or a field decl, so both must be handled.
-  virtual bool handleSyclAccessorType(const CXXRecordDecl *,
-                                      const CXXBaseSpecifier &, QualType) {
-    return true;
-  }
-  virtual bool handleSyclAccessorType(FieldDecl *, QualType) { return true; }
-  virtual bool handleSyclSamplerType(const CXXRecordDecl *,
+  // SYCL special class can be a base class or a field decl, so both must be
+  // handled.
+  virtual bool handleSyclSpecialType(const CXXRecordDecl *,
                                      const CXXBaseSpecifier &, QualType) {
     return true;
   }
-  virtual bool handleSyclSamplerType(FieldDecl *, QualType) { return true; }
+  virtual bool handleSyclSpecialType(FieldDecl *, QualType) { return true; }
+
   virtual bool handleSyclSpecConstantType(FieldDecl *, QualType) {
     return true;
   }
-  virtual bool handleSyclStreamType(const CXXRecordDecl *,
-                                    const CXXBaseSpecifier &, QualType) {
-    return true;
-  }
-  virtual bool handleSyclStreamType(FieldDecl *, QualType) { return true; }
+
   virtual bool handleSyclHalfType(const CXXRecordDecl *,
                                   const CXXBaseSpecifier &, QualType) {
     return true;
@@ -1671,10 +1646,9 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
     return false;
   }
 
-  bool checkAccessorType(QualType Ty, SourceRange Loc) {
-    assert(Util::isSyclAccessorType(Ty) &&
-           "Should only be called on SYCL accessor types.");
-
+  bool checkSyclSpecialType(QualType Ty, SourceRange Loc) {
+    assert(Util::isSyclSpecialType(Ty) &&
+           "Should only be called on sycl special class types.");
     const RecordDecl *RecD = Ty->getAsRecordDecl();
     if (const ClassTemplateSpecializationDecl *CTSD =
             dyn_cast<ClassTemplateSpecializationDecl>(RecD)) {
@@ -1706,14 +1680,14 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
     return isValid();
   }
 
-  bool handleSyclAccessorType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
-                              QualType FieldTy) final {
-    IsInvalid |= checkAccessorType(FieldTy, BS.getBeginLoc());
+  bool handleSyclSpecialType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
+                             QualType FieldTy) final {
+    IsInvalid |= checkSyclSpecialType(FieldTy, BS.getBeginLoc());
     return isValid();
   }
 
-  bool handleSyclAccessorType(FieldDecl *FD, QualType FieldTy) final {
-    IsInvalid |= checkAccessorType(FieldTy, FD->getLocation());
+  bool handleSyclSpecialType(FieldDecl *FD, QualType FieldTy) final {
+    IsInvalid |= checkSyclSpecialType(FieldTy, FD->getLocation());
     return isValid();
   }
 
@@ -1773,32 +1747,14 @@ class SyclKernelUnionChecker : public SyclKernelFieldHandler {
     return true;
   }
 
-  bool handleSyclAccessorType(FieldDecl *FD, QualType FieldTy) final {
-    return checkType(FD->getLocation(), FieldTy);
-  }
-
-  bool handleSyclAccessorType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
-                              QualType FieldTy) final {
-    return checkType(BS.getBeginLoc(), FieldTy);
-  }
-
-  bool handleSyclSamplerType(FieldDecl *FD, QualType FieldTy) final {
+  bool handleSyclSpecialType(FieldDecl *FD, QualType FieldTy) final {
     return checkType(FD->getLocation(), FieldTy);
   }
 
-  bool handleSyclSamplerType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
+  bool handleSyclSpecialType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
                              QualType FieldTy) final {
     return checkType(BS.getBeginLoc(), FieldTy);
   }
-
-  bool handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final {
-    return checkType(FD->getLocation(), FieldTy);
-  }
-
-  bool handleSyclStreamType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
-                            QualType FieldTy) final {
-    return checkType(BS.getBeginLoc(), FieldTy);
-  }
 };
 
 // A type to mark whether a collection requires decomposition.
@@ -1815,38 +1771,21 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler {
     CollectionStack.push_back(true);
   }
 
-  bool handleSyclAccessorType(const CXXRecordDecl *, const CXXBaseSpecifier &,
-                              QualType) final {
-    CollectionStack.back() = true;
-    return true;
-  }
-  bool handleSyclAccessorType(FieldDecl *, QualType) final {
-    CollectionStack.back() = true;
-    return true;
-  }
-
-  bool handleSyclSamplerType(const CXXRecordDecl *, const CXXBaseSpecifier &,
+  bool handleSyclSpecialType(const CXXRecordDecl *, const CXXBaseSpecifier &,
                              QualType) final {
     CollectionStack.back() = true;
     return true;
   }
-  bool handleSyclSamplerType(FieldDecl *, QualType) final {
+  bool handleSyclSpecialType(FieldDecl *, QualType) final {
     CollectionStack.back() = true;
     return true;
   }
+
   bool handleSyclSpecConstantType(FieldDecl *, QualType) final {
     CollectionStack.back() = true;
     return true;
   }
-  bool handleSyclStreamType(const CXXRecordDecl *, const CXXBaseSpecifier &,
-                            QualType) final {
-    CollectionStack.back() = true;
-    return true;
-  }
-  bool handleSyclStreamType(FieldDecl *, QualType) final {
-    CollectionStack.back() = true;
-    return true;
-  }
+
   bool handleSyclHalfType(const CXXRecordDecl *, const CXXBaseSpecifier &,
                           QualType) final {
     CollectionStack.back() = true;
@@ -2014,16 +1953,40 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
         SYCLIntelBufferLocationAttr::CreateImplicit(Ctx, LocationID));
   }
 
+  // Additional processing is required for accessor type.
+  void handleAccessorType(const CXXRecordDecl *RecordDecl, SourceLocation Loc) {
+    handleAccessorPropertyList(Params.back(), RecordDecl, Loc);
+    if (KernelDecl->hasAttr<SYCLSimdAttr>())
+      // In ESIMD, the kernels accessor's pointer argument needs to be marked.
+      Params.back()->addAttr(
+          SYCLSimdAccessorPtrAttr::CreateImplicit(SemaRef.getASTContext()));
+    // Get access mode of accessor.
+    const auto *AccessorSpecializationDecl =
+        cast<ClassTemplateSpecializationDecl>(RecordDecl);
+    const TemplateArgument &AccessModeArg =
+        AccessorSpecializationDecl->getTemplateArgs().get(2);
+
+    // Add implicit attribute to parameter decl when it is a read only
+    // SYCL accessor.
+    if (isReadOnlyAccessor(AccessModeArg))
+      Params.back()->addAttr(
+          SYCLAccessorReadonlyAttr::CreateImplicit(SemaRef.getASTContext()));
+  }
+
   // All special SYCL objects must have __init method. We extract types for
   // kernel parameters from __init method parameters. We will use __init method
   // and kernel parameters which we build here to initialize special objects in
-  // the kernel body. Accessors require additional processing and are handled in
-  // handleSyclAccessorType.
+  // the kernel body.
   bool handleSpecialType(FieldDecl *FD, QualType FieldTy) {
     const auto *RecordDecl = FieldTy->getAsCXXRecordDecl();
-    assert(RecordDecl && "The stream/sampler must be a RecordDecl");
-    CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName);
-    assert(InitMethod && "The stream/sampler must have the __init method");
+    assert(RecordDecl && "The type must be a RecordDecl");
+    llvm::StringLiteral MethodName =
+        KernelDecl->hasAttr<SYCLSimdAttr>() &&
+                Util::isSyclType(FieldTy, "accessor", true /*Tmp*/)
+            ? InitESIMDMethodName
+            : InitMethodName;
+    CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName);
+    assert(InitMethod && "The type must have the __init method");
 
     // Don't do -1 here because we count on this to be the first parameter added
     // (if any).
@@ -2031,6 +1994,15 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     for (const ParmVarDecl *Param : InitMethod->parameters()) {
       QualType ParamTy = Param->getType();
       addParam(FD, ParamTy.getCanonicalType());
+      // FIXME: This code is temporary, and will be removed once __init_esimd
+      // is removed and property list refactored.
+      // The function handleAccessorType includes a call to
+      // handleAccessorPropertyList. If new classes with property list are
+      // added, this code needs to be refactored to call
+      // handleAccessorPropertyList for each class which requires it.
+      if (ParamTy.getTypePtr()->isPointerType() &&
+          Util::isSyclType(FieldTy, "accessor", true /*Tmp*/))
+        handleAccessorType(RecordDecl, FD->getBeginLoc());
     }
     LastParamIndex = ParamIndex;
     return true;
@@ -2119,23 +2091,17 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     return true;
   }
 
-  // FIXME: Refactor accessor handling. There is a discrepancy in how
-  // inherited accessors are handled.
-  bool handleSyclAccessorType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
-                              QualType FieldTy) final {
+  bool handleSyclSpecialType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
+                             QualType FieldTy) final {
     const auto *RecordDecl = FieldTy->getAsCXXRecordDecl();
-    assert(RecordDecl && "The accessor/sampler must be a RecordDecl");
-    llvm::StringLiteral MethodName = KernelDecl->hasAttr<SYCLSimdAttr>()
-                                         ? InitESIMDMethodName
-                                         : InitMethodName;
+    assert(RecordDecl && "The type must be a RecordDecl");
+    llvm::StringLiteral MethodName =
+        KernelDecl->hasAttr<SYCLSimdAttr>() &&
+                Util::isSyclType(FieldTy, "accessor", true /*Tmp*/)
+            ? InitESIMDMethodName
+            : InitMethodName;
     CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName);
-    assert(InitMethod && "The accessor/sampler must have the __init method");
-
-    // Get access mode of accessor.
-    const auto *AccessorSpecializationDecl =
-        cast<ClassTemplateSpecializationDecl>(RecordDecl);
-    const TemplateArgument &AccessModeArg =
-        AccessorSpecializationDecl->getTemplateArgs().get(2);
+    assert(InitMethod && "The type must have the __init method");
 
     // Don't do -1 here because we count on this to be the first parameter added
     // (if any).
@@ -2143,62 +2109,21 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     for (const ParmVarDecl *Param : InitMethod->parameters()) {
       QualType ParamTy = Param->getType();
       addParam(BS, ParamTy.getCanonicalType());
-      if (ParamTy.getTypePtr()->isPointerType()) {
-        handleAccessorPropertyList(Params.back(), RecordDecl, BS.getBeginLoc());
-
-        // Add implicit attribute to parameter decl when it is a read only
-        // SYCL accessor.
-        if (isReadOnlyAccessor(AccessModeArg))
-          Params.back()->addAttr(SYCLAccessorReadonlyAttr::CreateImplicit(
-              SemaRef.getASTContext()));
-      }
+      // FIXME: This code is temporary, and will be removed once __init_esimd
+      // is removed and property list refactored.
+      // The function handleAccessorType includes a call to
+      // handleAccessorPropertyList. If new classes with property list are
+      // added, this code needs to be refactored to call
+      // handleAccessorPropertyList for each class which requires it.
+      if (ParamTy.getTypePtr()->isPointerType() &&
+          Util::isSyclType(FieldTy, "accessor", true /*Tmp*/))
+        handleAccessorType(RecordDecl, BS.getBeginLoc());
     }
     LastParamIndex = ParamIndex;
     return true;
   }
 
-  bool handleSyclAccessorType(FieldDecl *FD, QualType FieldTy) final {
-    const auto *RecordDecl = FieldTy->getAsCXXRecordDecl();
-    assert(RecordDecl && "The accessor must be a RecordDecl");
-
-    // Get access mode of accessor.
-    const auto *AccessorSpecializationDecl =
-        cast<ClassTemplateSpecializationDecl>(RecordDecl);
-    const TemplateArgument &AccessModeArg =
-        AccessorSpecializationDecl->getTemplateArgs().get(2);
-
-    llvm::StringLiteral MethodName = KernelDecl->hasAttr<SYCLSimdAttr>()
-                                         ? InitESIMDMethodName
-                                         : InitMethodName;
-    CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName);
-    assert(InitMethod && "The accessor must have the __init method");
-
-    // Don't do -1 here because we count on this to be the first parameter added
-    // (if any).
-    size_t ParamIndex = Params.size();
-    for (const ParmVarDecl *Param : InitMethod->parameters()) {
-      QualType ParamTy = Param->getType();
-      addParam(FD, ParamTy.getCanonicalType());
-      if (ParamTy.getTypePtr()->isPointerType()) {
-        handleAccessorPropertyList(Params.back(), RecordDecl,
-                                   FD->getLocation());
-        if (KernelDecl->hasAttr<SYCLSimdAttr>())
-          // In ESIMD kernels accessor's pointer argument needs to be marked
-          Params.back()->addAttr(
-              SYCLSimdAccessorPtrAttr::CreateImplicit(SemaRef.getASTContext()));
-
-        // Add implicit attribute to parameter decl when it is a read only
-        // SYCL accessor.
-        if (isReadOnlyAccessor(AccessModeArg))
-          Params.back()->addAttr(SYCLAccessorReadonlyAttr::CreateImplicit(
-              SemaRef.getASTContext()));
-      }
-    }
-    LastParamIndex = ParamIndex;
-    return true;
-  }
-
-  bool handleSyclSamplerType(FieldDecl *FD, QualType FieldTy) final {
+  bool handleSyclSpecialType(FieldDecl *FD, QualType FieldTy) final {
     return handleSpecialType(FD, FieldTy);
   }
 
@@ -2283,17 +2208,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     return true;
   }
 
-  bool handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final {
-    return handleSpecialType(FD, FieldTy);
-  }
-
-  bool handleSyclStreamType(const CXXRecordDecl *, const CXXBaseSpecifier &,
-                            QualType FieldTy) final {
-    // FIXME SYCL stream should be usable as a base type
-    // See https://github.com/intel/llvm/issues/1552
-    return true;
-  }
-
   // Generate kernel argument to intialize specialization constants. This
   // argument is only generated when the target has no native support for
   // specialization constants
@@ -2316,7 +2230,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
                                    std::end(Params));
   }
   using SyclKernelFieldHandler::handleSyclHalfType;
-  using SyclKernelFieldHandler::handleSyclSamplerType;
 };
 
 class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler {
@@ -2331,11 +2244,11 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler {
 
   bool handleSpecialType(QualType FieldTy) {
     const CXXRecordDecl *RecordDecl = FieldTy->getAsCXXRecordDecl();
-    assert(RecordDecl && "The accessor/sampler must be a RecordDecl");
+    assert(RecordDecl && "The type must be a RecordDecl");
     llvm::StringLiteral MethodName =
         IsSIMD ? InitESIMDMethodName : InitMethodName;
     CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName);
-    assert(InitMethod && "The accessor/sampler must have the __init method");
+    assert(InitMethod && "The type must have the __init method");
     for (const ParmVarDecl *Param : InitMethod->parameters())
       addParam(Param->getType());
     return true;
@@ -2352,20 +2265,11 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler {
           << SizeOfParams << MaxKernelArgsSize;
   }
 
-  bool handleSyclAccessorType(FieldDecl *FD, QualType FieldTy) final {
-    return handleSpecialType(FieldTy);
-  }
-
-  bool handleSyclAccessorType(const CXXRecordDecl *, const CXXBaseSpecifier &,
-                              QualType FieldTy) final {
+  bool handleSyclSpecialType(FieldDecl *FD, QualType FieldTy) final {
     return handleSpecialType(FieldTy);
   }
 
-  bool handleSyclSamplerType(FieldDecl *FD, QualType FieldTy) final {
-    return handleSpecialType(FieldTy);
-  }
-
-  bool handleSyclSamplerType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
+  bool handleSyclSpecialType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
                              QualType FieldTy) final {
     return handleSpecialType(FieldTy);
   }
@@ -2405,57 +2309,14 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler {
     addParam(FieldTy);
     return true;
   }
-
-  bool handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final {
-    addParam(FieldTy);
-    return true;
-  }
-  bool handleSyclStreamType(const CXXRecordDecl *, const CXXBaseSpecifier &,
-                            QualType FieldTy) final {
-    addParam(FieldTy);
-    return true;
-  }
   using SyclKernelFieldHandler::handleSyclHalfType;
 };
 
-enum class KernelArgDescription {
-  BaseClass,
-  DecomposedMember,
-  WrappedPointer,
-  WrappedArray,
-  Accessor,
-  AccessorBase,
-  Sampler,
-  Stream,
-  KernelHandler,
-  None
-};
-
-StringRef getKernelArgDesc(KernelArgDescription Desc) {
-  switch (Desc) {
-  case KernelArgDescription::BaseClass:
-    return "Compiler generated argument for base class,";
-  case KernelArgDescription::DecomposedMember:
-    return "Compiler generated argument for decomposed struct/class,";
-  case KernelArgDescription::WrappedPointer:
-    return "Compiler generated argument for nested pointer,";
-  case KernelArgDescription::WrappedArray:
-    return "Compiler generated argument for array,";
-  case KernelArgDescription::Accessor:
-    return "Compiler generated argument for accessor,";
-  case KernelArgDescription::AccessorBase:
-    return "Compiler generated argument for accessor base class,";
-  case KernelArgDescription::Sampler:
-    return "Compiler generated argument for sampler,";
-  case KernelArgDescription::Stream:
-    return "Compiler generated argument for stream,";
-  case KernelArgDescription::KernelHandler:
-    return "Compiler generated argument for SYCL2020 specialization constant";
-  case KernelArgDescription::None:
+std::string getKernelArgDesc(StringRef KernelArgDescription) {
+  if (KernelArgDescription == ":" || KernelArgDescription == "")
     return "";
-  }
-  llvm_unreachable(
-      "switch should cover all possible values for KernelArgDescription");
+  return ("Compiler generated argument for " + KernelArgDescription + ",")
+      .str();
 }
 
 class SyclOptReportCreator : public SyclKernelFieldHandler {
@@ -2463,17 +2324,20 @@ class SyclOptReportCreator : public SyclKernelFieldHandler {
   SourceLocation KernelInvocationLoc;
 
   void addParam(const FieldDecl *KernelArg, QualType KernelArgType,
-                KernelArgDescription KernelArgDesc) {
+                StringRef KernelArgDescription) {
     StringRef NameToEmitInDescription = KernelArg->getName();
     const RecordDecl *KernelArgParent = KernelArg->getParent();
-    if (KernelArgParent &&
-        KernelArgDesc == KernelArgDescription::DecomposedMember) {
+    if (KernelArgParent && KernelArgDescription == "decomposed struct/class")
       NameToEmitInDescription = KernelArgParent->getName();
-    }
 
-    bool isWrappedField =
-        KernelArgDesc == KernelArgDescription::WrappedPointer ||
-        KernelArgDesc == KernelArgDescription::WrappedArray;
+    bool isWrappedField = KernelArgDescription == "WrappedPointer" ||
+                          KernelArgDescription == "WrappedArray";
+
+    KernelArgDescription =
+        (KernelArgDescription == "WrappedPointer"
+             ? "nested pointer"
+             : (KernelArgDescription == "WrappedArray" ? "array"
+                                                       : KernelArgDescription));
 
     unsigned KernelArgSize =
         SemaRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity();
@@ -2481,49 +2345,43 @@ class SyclOptReportCreator : public SyclKernelFieldHandler {
     SemaRef.getDiagnostics().getSYCLOptReport().AddKernelArgs(
         DC.getKernelDecl(), NameToEmitInDescription,
         isWrappedField ? "Compiler generated" : KernelArgType.getAsString(),
-        KernelInvocationLoc, KernelArgSize, getKernelArgDesc(KernelArgDesc),
-        (KernelArgDesc == KernelArgDescription::DecomposedMember)
+        KernelInvocationLoc, KernelArgSize,
+        getKernelArgDesc(KernelArgDescription),
+        (KernelArgDescription == "decomposed struct/class")
             ? ("Field:" + KernelArg->getName().str() + ", ")
             : "");
   }
 
   void addParam(const FieldDecl *FD, QualType FieldTy) {
-    KernelArgDescription Desc = KernelArgDescription::None;
+    std::string KernelArgDescription = FieldTy.getAsString();
     const RecordDecl *RD = FD->getParent();
+    if (FieldTy->isScalarType())
+      KernelArgDescription = "";
     if (RD && RD->hasAttr<SYCLRequiresDecompositionAttr>())
-      Desc = KernelArgDescription::DecomposedMember;
+      KernelArgDescription = "decomposed struct/class";
 
-    addParam(FD, FieldTy, Desc);
+    addParam(FD, FieldTy, KernelArgDescription);
   }
 
   // Handles base classes.
   void addParam(const CXXBaseSpecifier &, QualType KernelArgType,
-                KernelArgDescription KernelArgDesc) {
+                StringRef KernelArgDescription) {
     unsigned KernelArgSize =
         SemaRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity();
     SemaRef.getDiagnostics().getSYCLOptReport().AddKernelArgs(
         DC.getKernelDecl(), KernelArgType.getAsString(),
         KernelArgType.getAsString(), KernelInvocationLoc, KernelArgSize,
-        getKernelArgDesc(KernelArgDesc), "");
+        getKernelArgDesc(KernelArgDescription), "");
   }
 
   // Handles specialization constants.
-  void addParam(QualType KernelArgType, KernelArgDescription KernelArgDesc) {
+  void addParam(QualType KernelArgType, std::string KernelArgDescription) {
     unsigned KernelArgSize =
         SemaRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity();
     SemaRef.getDiagnostics().getSYCLOptReport().AddKernelArgs(
         DC.getKernelDecl(), "", KernelArgType.getAsString(),
-        KernelInvocationLoc, KernelArgSize, getKernelArgDesc(KernelArgDesc),
-        "");
-  }
-
-  // Handles SYCL special types (accessor, sampler and stream) and modified
-  // types (arrays and pointers)
-  bool handleSpecialType(const FieldDecl *FD, QualType FieldTy,
-                         KernelArgDescription Desc) {
-    for (const auto *Param : DC.getParamVarDeclsForCurrentField())
-      addParam(FD, Param->getType(), Desc);
-    return true;
+        KernelInvocationLoc, KernelArgSize,
+        getKernelArgDesc(KernelArgDescription), "");
   }
 
 public:
@@ -2531,13 +2389,15 @@ class SyclOptReportCreator : public SyclKernelFieldHandler {
   SyclOptReportCreator(Sema &S, SyclKernelDeclCreator &DC, SourceLocation Loc)
       : SyclKernelFieldHandler(S), DC(DC), KernelInvocationLoc(Loc) {}
 
-  bool handleSyclAccessorType(FieldDecl *FD, QualType FieldTy) final {
-    return handleSpecialType(
-        FD, FieldTy, KernelArgDescription(KernelArgDescription::Accessor));
+  bool handleSyclSpecialType(FieldDecl *FD, QualType FieldTy) final {
+    for (const auto *Param : DC.getParamVarDeclsForCurrentField())
+      addParam(FD, Param->getType(), FieldTy.getAsString());
+    return true;
   }
 
-  bool handleSyclAccessorType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
-                              QualType FieldTy) final {
+  bool handleSyclSpecialType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
+                             QualType FieldTy) final {
+    std::string KernelArgDescription = "base class " + FieldTy.getAsString();
     for (const auto *Param : DC.getParamVarDeclsForCurrentField()) {
       QualType KernelArgType = Param->getType();
       unsigned KernelArgSize = SemaRef.getASTContext()
@@ -2546,27 +2406,22 @@ class SyclOptReportCreator : public SyclKernelFieldHandler {
       SemaRef.getDiagnostics().getSYCLOptReport().AddKernelArgs(
           DC.getKernelDecl(), FieldTy.getAsString(),
           KernelArgType.getAsString(), KernelInvocationLoc, KernelArgSize,
-          getKernelArgDesc(
-              KernelArgDescription(KernelArgDescription::AccessorBase)),
-          "");
+          getKernelArgDesc(KernelArgDescription), "");
     }
     return true;
   }
 
-  bool handleSyclSamplerType(FieldDecl *FD, QualType FieldTy) final {
-    return handleSpecialType(
-        FD, FieldTy, KernelArgDescription(KernelArgDescription::Sampler));
-  }
-
   bool handlePointerType(FieldDecl *FD, QualType FieldTy) final {
-    KernelArgDescription Desc = KernelArgDescription::None;
+    std::string KernelArgDescription = ":";
     ParmVarDecl *KernelParameter = DC.getParamVarDeclsForCurrentField()[0];
     // Compiler generated openCL kernel argument for current pointer field
     // is not a pointer. This means we are processing a nested pointer and
     // the openCL kernel argument is of type __wrapper_class.
     if (!KernelParameter->getType()->isPointerType())
-      Desc = KernelArgDescription::WrappedPointer;
-    return handleSpecialType(FD, FieldTy, Desc);
+      KernelArgDescription = "WrappedPointer";
+    for (const auto *Param : DC.getParamVarDeclsForCurrentField())
+      addParam(FD, Param->getType(), KernelArgDescription);
+    return true;
   }
 
   bool handleScalarType(FieldDecl *FD, QualType FieldTy) final {
@@ -2576,8 +2431,8 @@ class SyclOptReportCreator : public SyclKernelFieldHandler {
 
   bool handleSimpleArrayType(FieldDecl *FD, QualType FieldTy) final {
     // Simple arrays are always wrapped.
-    handleSpecialType(FD, FieldTy,
-                      KernelArgDescription(KernelArgDescription::WrappedArray));
+    for (const auto *Param : DC.getParamVarDeclsForCurrentField())
+      addParam(FD, Param->getType(), "WrappedArray");
     return true;
   }
 
@@ -2589,7 +2444,7 @@ class SyclOptReportCreator : public SyclKernelFieldHandler {
 
   bool handleNonDecompStruct(const CXXRecordDecl *Base,
                              const CXXBaseSpecifier &BS, QualType Ty) final {
-    addParam(BS, Ty, KernelArgDescription(KernelArgDescription::BaseClass));
+    addParam(BS, Ty, "base class");
     return true;
   }
 
@@ -2602,21 +2457,14 @@ class SyclOptReportCreator : public SyclKernelFieldHandler {
     return true;
   }
 
-  bool handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final {
-    return handleSpecialType(
-        FD, FieldTy, KernelArgDescription(KernelArgDescription::Stream));
-  }
-
   void handleSyclKernelHandlerType() {
     ASTContext &Context = SemaRef.getASTContext();
     if (isDefaultSPIRArch(Context))
       return;
     addParam(DC.getParamVarDeclsForCurrentField()[0]->getType(),
-             KernelArgDescription(KernelArgDescription::KernelHandler));
+             "SYCL2020 specialization constant");
   }
   using SyclKernelFieldHandler::handleSyclHalfType;
-  using SyclKernelFieldHandler::handleSyclSamplerType;
-  using SyclKernelFieldHandler::handleSyclStreamType;
 };
 
 static CXXMethodDecl *getOperatorParens(const CXXRecordDecl *Rec) {
@@ -2980,10 +2828,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
     createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts);
     CXXMethodDecl *FinalizeMethod =
         getMethodByName(RecordDecl, FinalizeMethodName);
-    // A finalize-method is expected for stream class.
-    if (!FinalizeMethod && Util::isSyclStreamType(Ty))
-      SemaRef.Diag(FD->getLocation(), diag::err_sycl_expected_finalize_method);
-    else
+    // A finalize-method is expected for special type such as stream.
+    if (FinalizeMethod)
       createSpecialMethodCall(RecordDecl, FinalizeMethodName, FinalizeStmts);
 
     removeFieldMemberExpr(FD, Ty);
@@ -3060,34 +2906,19 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
     DeclCreator.setBody(KernelBody);
   }
 
-  bool handleSyclAccessorType(FieldDecl *FD, QualType Ty) final {
+  bool handleSyclSpecialType(FieldDecl *FD, QualType Ty) final {
     return handleSpecialType(FD, Ty);
   }
 
-  bool handleSyclAccessorType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
-                              QualType Ty) final {
+  bool handleSyclSpecialType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
+                             QualType Ty) final {
     return handleSpecialType(BS, Ty);
   }
 
-  bool handleSyclSamplerType(FieldDecl *FD, QualType Ty) final {
-    return handleSpecialType(FD, Ty);
-  }
-
   bool handleSyclSpecConstantType(FieldDecl *FD, QualType Ty) final {
     return handleSpecialType(FD, Ty);
   }
 
-  bool handleSyclStreamType(FieldDecl *FD, QualType Ty) final {
-    return handleSpecialType(FD, Ty);
-  }
-
-  bool handleSyclStreamType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
-                            QualType Ty) final {
-    // FIXME SYCL stream should be usable as a base type
-    // See https://github.com/intel/llvm/issues/1552
-    return true;
-  }
-
   bool handleSyclHalfType(FieldDecl *FD, QualType Ty) final {
     addSimpleFieldInit(FD, Ty);
     return true;
@@ -3256,7 +3087,6 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
   }
 
   using SyclKernelFieldHandler::handleSyclHalfType;
-  using SyclKernelFieldHandler::handleSyclSamplerType;
 };
 
 // Kernels are only the unnamed-lambda feature if the feature is enabled, AND
@@ -3375,9 +3205,9 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
     setThisItemIsCalled(KernelFunc);
   }
 
-  bool handleSyclAccessorType(const CXXRecordDecl *RD,
-                              const CXXBaseSpecifier &BC,
-                              QualType FieldTy) final {
+  bool handleSyclSpecialType(const CXXRecordDecl *RD,
+                             const CXXBaseSpecifier &BC,
+                             QualType FieldTy) final {
     const auto *AccTy =
         cast<ClassTemplateSpecializationDecl>(FieldTy->getAsRecordDecl());
     assert(AccTy->getTemplateArgs().size() >= 2 &&
@@ -3391,37 +3221,32 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
     return true;
   }
 
-  bool handleSyclAccessorType(FieldDecl *FD, QualType FieldTy) final {
-    const auto *AccTy =
-        cast<ClassTemplateSpecializationDecl>(FieldTy->getAsRecordDecl());
-    assert(AccTy->getTemplateArgs().size() >= 2 &&
-           "Incorrect template args for Accessor Type");
-    int Dims = static_cast<int>(
-        AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue());
-    int Info = getAccessTarget(AccTy) | (Dims << 11);
-
-    Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info,
-                        CurOffset + offsetOf(FD, FieldTy));
-    return true;
-  }
-
-  bool handleSyclSamplerType(FieldDecl *FD, QualType FieldTy) final {
-    const auto *SamplerTy = FieldTy->getAsCXXRecordDecl();
-    assert(SamplerTy && "Sampler type must be a C++ record type");
-    CXXMethodDecl *InitMethod = getMethodByName(SamplerTy, InitMethodName);
-    assert(InitMethod && "sampler must have __init method");
-
-    // sampler __init method has only one argument
-    const ParmVarDecl *SamplerArg = InitMethod->getParamDecl(0);
-    assert(SamplerArg && "sampler __init method must have sampler parameter");
-
-    // For samplers, we do some special work to ONLY initialize the first item
-    // to the InitMethod as a performance improvement presumably, so the normal
-    // offsetOf calculation wouldn't work correctly. Therefore, we need to call
-    // a version of addParam where we calculate the offset based on the true
-    // FieldDecl/FieldType pair, rather than the SampleArg type.
-    addParam(SamplerArg->getType(), SYCLIntegrationHeader::kind_sampler,
-             offsetOf(FD, FieldTy));
+  bool handleSyclSpecialType(FieldDecl *FD, QualType FieldTy) final {
+    const auto *ClassTy = FieldTy->getAsCXXRecordDecl();
+    assert(ClassTy && "Type must be a C++ record type");
+    if (Util::isSyclType(FieldTy, "accessor", true /*Tmp*/)) {
+      const auto *AccTy =
+          dyn_cast<ClassTemplateSpecializationDecl>(FieldTy->getAsRecordDecl());
+      assert(AccTy->getTemplateArgs().size() >= 2 &&
+             "Incorrect template args for Accessor Type");
+      int Dims = static_cast<int>(
+          AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue());
+      int Info = getAccessTarget(AccTy) | (Dims << 11);
+
+      Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info,
+                          CurOffset + offsetOf(FD, FieldTy));
+    } else {
+      if (getMethodByName(ClassTy, FinalizeMethodName))
+        addParam(FD, FieldTy, SYCLIntegrationHeader::kind_stream);
+      else {
+        CXXMethodDecl *InitMethod = getMethodByName(ClassTy, InitMethodName);
+        assert(InitMethod && "type must have __init method");
+        const ParmVarDecl *SamplerArg = InitMethod->getParamDecl(0);
+        assert(SamplerArg && "Init method must have arguments");
+        addParam(SamplerArg->getType(), SYCLIntegrationHeader::kind_sampler,
+                 offsetOf(FD, FieldTy));
+      }
+    }
     return true;
   }
 
@@ -3476,18 +3301,6 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
     return handleScalarType(FD, FieldTy);
   }
 
-  bool handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final {
-    addParam(FD, FieldTy, SYCLIntegrationHeader::kind_stream);
-    return true;
-  }
-
-  bool handleSyclStreamType(const CXXRecordDecl *, const CXXBaseSpecifier &BC,
-                            QualType FieldTy) final {
-    // FIXME SYCL stream should be usable as a base type
-    // See https://github.com/intel/llvm/issues/1552
-    return true;
-  }
-
   bool handleSyclHalfType(FieldDecl *FD, QualType FieldTy) final {
     addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout);
     return true;
@@ -3552,7 +3365,6 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
 
   using SyclKernelFieldHandler::enterStruct;
   using SyclKernelFieldHandler::handleSyclHalfType;
-  using SyclKernelFieldHandler::handleSyclSamplerType;
   using SyclKernelFieldHandler::leaveStruct;
 };
 
@@ -5190,15 +5002,13 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
 // -----------------------------------------------------------------------------
 // Utility class methods
 // -----------------------------------------------------------------------------
-
-bool Util::isSyclAccessorType(QualType Ty) {
-  return isSyclType(Ty, "accessor", true /*Tmpl*/);
+bool Util::isSyclSpecialType(const QualType Ty) {
+  const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl();
+  if (!RecTy)
+    return false;
+  return RecTy->hasAttr<SYCLSpecialClassAttr>();
 }
 
-bool Util::isSyclSamplerType(QualType Ty) { return isSyclType(Ty, "sampler"); }
-
-bool Util::isSyclStreamType(QualType Ty) { return isSyclType(Ty, "stream"); }
-
 bool Util::isSyclHalfType(QualType Ty) {
   std::array<DeclContextDesc, 5> Scopes = {
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
index cd0dfa8b9b71a..8cad0b9c082a8 100644
--- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -13,7 +13,7 @@ struct sampler_impl {
 #endif
 };
 
-class sampler {
+class __attribute__((sycl_special_class)) sampler {
   struct sampler_impl impl;
 #ifdef __SYCL_DEVICE_ONLY__
   void __init(__ocl_sampler_t Sampler) { impl.m_Sampler = Sampler; }
@@ -184,7 +184,7 @@ template <typename dataT, int dimensions, access::mode accessmode,
           access::target accessTarget = access::target::global_buffer,
           access::placeholder isPlaceholder = access::placeholder::false_t,
           typename propertyListT = ext::oneapi::accessor_property_list<>>
-class accessor {
+class __attribute__((sycl_special_class)) accessor {
 
 public:
   void use(void) const {}
@@ -250,7 +250,7 @@ struct _ImageImplT {
 };
 
 template <typename dataT, int dimensions, access::mode accessmode>
-class accessor<dataT, dimensions, accessmode, access::target::image, access::placeholder::false_t> {
+class __attribute__((sycl_special_class)) accessor<dataT, dimensions, accessmode, access::target::image, access::placeholder::false_t> {
 public:
   void use(void) const {}
   template <typename... T>
@@ -447,7 +447,7 @@ class handler {
   }
 };
 
-class stream {
+class __attribute__((sycl_special_class)) stream {
 public:
   stream(unsigned long BufferSize, unsigned long MaxStatementSize,
          handler &CGH) {}
diff --git a/clang/test/CodeGenSYCL/accessor-readonly-invalid-lib.cpp b/clang/test/CodeGenSYCL/accessor-readonly-invalid-lib.cpp
index a085f3da0bcce..ba334aaad8abe 100644
--- a/clang/test/CodeGenSYCL/accessor-readonly-invalid-lib.cpp
+++ b/clang/test/CodeGenSYCL/accessor-readonly-invalid-lib.cpp
@@ -60,7 +60,7 @@ struct _ImplT {
 template <typename dataT, int dimensions, access::mode accessmode,
           access::target accessTarget = access::target::global_buffer,
           access::placeholder isPlaceholder = access::placeholder::false_t>
-class accessor {
+class __attribute__((sycl_special_class)) accessor {
 
 public:
   void use(void) const {}
diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
index 5d7a002925d8c..5cfc6a1518a51 100644
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -171,6 +171,7 @@
 // CHECK-NEXT: SYCLIntelUseStallEnableClusters (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLRegisterNum (SubjectMatchRule_variable_is_global)
 // CHECK-NEXT: SYCLSimd (SubjectMatchRule_function, SubjectMatchRule_variable_is_global)
+// CHECK-NEXT: SYCLSpecialClass (SubjectMatchRule_record)
 // CHECK-NEXT: ScopedLockable (SubjectMatchRule_record)
 // CHECK-NEXT: Section (SubjectMatchRule_function, SubjectMatchRule_variable_is_global, SubjectMatchRule_objc_method, SubjectMatchRule_objc_property)
 // CHECK-NEXT: SetTypestate (SubjectMatchRule_function_is_member)
diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp
index 73eaad209889b..90c5f3042cb66 100644
--- a/clang/test/SemaSYCL/Inputs/sycl.hpp
+++ b/clang/test/SemaSYCL/Inputs/sycl.hpp
@@ -107,8 +107,7 @@ template <typename dataT, int dimensions, access::mode accessmode,
           access::target accessTarget = access::target::global_buffer,
           access::placeholder isPlaceholder = access::placeholder::false_t,
           typename propertyListT = ext::oneapi::accessor_property_list<>>
-class accessor {
-
+class __attribute__((sycl_special_class)) accessor {
 public:
   void use(void) const {}
   void use(void *) const {}
@@ -167,7 +166,7 @@ struct _ImageImplT {
 };
 
 template <typename dataT, int dimensions, access::mode accessmode>
-class accessor<dataT, dimensions, accessmode, access::target::image, access::placeholder::false_t> {
+class __attribute__((sycl_special_class)) accessor<dataT, dimensions, accessmode, access::target::image, access::placeholder::false_t> {
 public:
   void use(void) const {}
   template <typename... T>
@@ -186,7 +185,7 @@ struct sampler_impl {
 #endif
 };
 
-class sampler {
+class __attribute__((sycl_special_class)) sampler {
   struct sampler_impl impl;
 #ifdef __SYCL_DEVICE_ONLY__
   void __init(__ocl_sampler_t Sampler) { impl.m_Sampler = Sampler; }
@@ -290,7 +289,7 @@ class handler {
   }
 };
 
-class stream {
+class __attribute__((sycl_special_class)) stream {
   accessor<int, 1, access::mode::read> acc;
 
 public:
diff --git a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp
index 384b4f91b6b2d..624f17900ac8a 100644
--- a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp
+++ b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp
@@ -132,7 +132,7 @@ int main() {
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '3'
 // SPIR-NEXT: String:          ':'
-// SPIR-NEXT: String:          Compiler generated argument for accessor,
+// SPIR-NEXT: String:          'Compiler generated argument for sycl::accessor<char, 1, sycl::access::mode::read>,'
 // SPIR-NEXT: String:          decompAcc
 // SPIR-NEXT: String:          '  ('
 // SPIR-NEXT: String:          ''
@@ -153,7 +153,7 @@ int main() {
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '4'
 // SPIR-NEXT: String:          ':'
-// SPIR-NEXT: String:          Compiler generated argument for accessor,
+// SPIR-NEXT: String:          'Compiler generated argument for sycl::accessor<char, 1, sycl::access::mode::read>,'
 // SPIR-NEXT: String:          decompAcc
 // SPIR-NEXT: String:          '  ('
 // SPIR-NEXT: String:          ''
@@ -174,7 +174,7 @@ int main() {
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '5'
 // SPIR-NEXT: String:          ':'
-// SPIR-NEXT: String:          Compiler generated argument for accessor,
+// SPIR-NEXT: String:          'Compiler generated argument for sycl::accessor<char, 1, sycl::access::mode::read>,'
 // SPIR-NEXT: String:          decompAcc
 // SPIR-NEXT: String:          '  ('
 // SPIR-NEXT: String:          ''
@@ -195,7 +195,7 @@ int main() {
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '6'
 // SPIR-NEXT: String:          ':'
-// SPIR-NEXT: String:          Compiler generated argument for accessor,
+// SPIR-NEXT: String:          'Compiler generated argument for sycl::accessor<char, 1, sycl::access::mode::read>,'
 // SPIR-NEXT: String:          decompAcc
 // SPIR-NEXT: String:          '  ('
 // SPIR-NEXT: String:          ''
@@ -216,7 +216,7 @@ int main() {
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '7'
 // SPIR-NEXT: String:          ':'
-// SPIR-NEXT: String:          Compiler generated argument for stream,
+// SPIR-NEXT: String:          'Compiler generated argument for sycl::stream,'
 // SPIR-NEXT: String:          DecompStream
 // SPIR-NEXT: String:          '  ('
 // SPIR-NEXT: String:          ''
@@ -237,7 +237,7 @@ int main() {
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '8'
 // SPIR-NEXT: String:          ':'
-// SPIR-NEXT: String:          Compiler generated argument for stream,
+// SPIR-NEXT: String:          'Compiler generated argument for sycl::stream,'
 // SPIR-NEXT: String:          DecompStream
 // SPIR-NEXT: String:          '  ('
 // SPIR-NEXT: String:          ''
@@ -258,7 +258,7 @@ int main() {
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '9'
 // SPIR-NEXT: String:          ':'
-// SPIR-NEXT: String:          Compiler generated argument for stream,
+// SPIR-NEXT: String:          'Compiler generated argument for sycl::stream,'
 // SPIR-NEXT: String:          DecompStream
 // SPIR-NEXT: String:          '  ('
 // SPIR-NEXT: String:          ''
@@ -279,7 +279,7 @@ int main() {
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '10'
 // SPIR-NEXT: String:          ':'
-// SPIR-NEXT: String:          Compiler generated argument for stream,
+// SPIR-NEXT: String:          'Compiler generated argument for sycl::stream,'
 // SPIR-NEXT: String:          DecompStream
 // SPIR-NEXT: String:          '  ('
 // SPIR-NEXT: String:          ''
@@ -300,7 +300,7 @@ int main() {
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '11'
 // SPIR-NEXT: String:          ':'
-// SPIR-NEXT: String:          Compiler generated argument for stream,
+// SPIR-NEXT: String:          'Compiler generated argument for sycl::stream,'
 // SPIR-NEXT: String:          DecompStream
 // SPIR-NEXT: String:          '  ('
 // SPIR-NEXT: String:          ''
@@ -384,7 +384,7 @@ int main() {
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '15'
 // SPIR-NEXT: String:          ':'
-// SPIR-NEXT: String:          Compiler generated argument for sampler,
+// SPIR-NEXT: String:          'Compiler generated argument for sycl::sampler,'
 // SPIR-NEXT: String:          Sampl
 // SPIR-NEXT: String:          '  ('
 // SPIR-NEXT: String:          ''
@@ -407,7 +407,7 @@ int main() {
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '0'
 // SPIR-NEXT: String:          ':'
-// SPIR-NEXT: String:          Compiler generated argument for accessor base class,
+// SPIR-NEXT: String:          'Compiler generated argument for base class sycl::accessor<char, 1, sycl::access::mode::read>,'
 // SPIR-NEXT: String:          'sycl::accessor<char, 1, sycl::access::mode::read>'
 // SPIR-NEXT: String:          '  ('
 // SPIR-NEXT: String:          ''
@@ -428,7 +428,7 @@ int main() {
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '1'
 // SPIR-NEXT: String:          ':'
-// SPIR-NEXT: String:          Compiler generated argument for accessor base class,
+// SPIR-NEXT: String:          'Compiler generated argument for base class sycl::accessor<char, 1, sycl::access::mode::read>,'
 // SPIR-NEXT: String:          'sycl::accessor<char, 1, sycl::access::mode::read>'
 // SPIR-NEXT: String:          '  ('
 // SPIR-NEXT: String:          ''
@@ -449,7 +449,7 @@ int main() {
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '2'
 // SPIR-NEXT: String:          ':'
-// SPIR-NEXT: String:          Compiler generated argument for accessor base class,
+// SPIR-NEXT: String:          'Compiler generated argument for base class sycl::accessor<char, 1, sycl::access::mode::read>,'
 // SPIR-NEXT: String:          'sycl::accessor<char, 1, sycl::access::mode::read>'
 // SPIR-NEXT: String:          '  ('
 // SPIR-NEXT: String:          ''
@@ -470,7 +470,7 @@ int main() {
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '3'
 // SPIR-NEXT: String:          ':'
-// SPIR-NEXT: String:          Compiler generated argument for accessor base class,
+// SPIR-NEXT: String:          'Compiler generated argument for base class sycl::accessor<char, 1, sycl::access::mode::read>,'
 // SPIR-NEXT: String:          'sycl::accessor<char, 1, sycl::access::mode::read>'
 // SPIR-NEXT: String:          '  ('
 // SPIR-NEXT: String:          ''
diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
old mode 100755
new mode 100644
index 5430b465a849d..012a257bb6707
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -784,7 +784,7 @@ class __image_array_slice__ {
 template <typename DataT, int Dimensions, access::mode AccessMode,
           access::target AccessTarget, access::placeholder IsPlaceholder,
           typename PropertyListT>
-class accessor :
+class __SYCL_SPECIAL_CLASS accessor :
 #ifndef __SYCL_DEVICE_ONLY__
     public detail::AccessorBaseHost,
 #endif
@@ -1578,7 +1578,7 @@ class accessor :
 
   template <int Dims = Dimensions>
   typename detail::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
-                               atomic<DataT, AS>>
+                             atomic<DataT, AS>>
   operator[](id<Dimensions> Index) const {
     const size_t LinearIndex = getLinearIndex(Index);
     return atomic<DataT, AS>(
@@ -1593,7 +1593,6 @@ class accessor :
     return atomic<DataT, AS>(
         multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
   }
-
   template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 1)>>
   typename AccessorCommonT::template AccessorSubscript<Dims - 1>
   operator[](size_t Index) const {
@@ -1798,8 +1797,8 @@ accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3,
 /// \ingroup sycl_api_acc
 template <typename DataT, int Dimensions, access::mode AccessMode,
           access::placeholder IsPlaceholder>
-class accessor<DataT, Dimensions, AccessMode, access::target::local,
-               IsPlaceholder> :
+class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
+                                    access::target::local, IsPlaceholder> :
 #ifndef __SYCL_DEVICE_ONLY__
     public detail::LocalAccessorBaseHost,
 #endif
@@ -1993,8 +1992,8 @@ class accessor<DataT, Dimensions, AccessMode, access::target::local,
 /// \ingroup sycl_api_acc
 template <typename DataT, int Dimensions, access::mode AccessMode,
           access::placeholder IsPlaceholder>
-class accessor<DataT, Dimensions, AccessMode, access::target::image,
-               IsPlaceholder>
+class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
+                                    access::target::image, IsPlaceholder>
     : public detail::image_accessor<DataT, Dimensions, AccessMode,
                                     access::target::image, IsPlaceholder> {
 public:
@@ -2052,8 +2051,8 @@ class accessor<DataT, Dimensions, AccessMode, access::target::image,
 /// \ingroup sycl_api_acc
 template <typename DataT, int Dimensions, access::mode AccessMode,
           access::placeholder IsPlaceholder>
-class accessor<DataT, Dimensions, AccessMode, access::target::host_image,
-               IsPlaceholder>
+class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
+                                    access::target::host_image, IsPlaceholder>
     : public detail::image_accessor<DataT, Dimensions, AccessMode,
                                     access::target::host_image, IsPlaceholder> {
 public:
@@ -2083,8 +2082,8 @@ class accessor<DataT, Dimensions, AccessMode, access::target::host_image,
 /// \ingroup sycl_api_acc
 template <typename DataT, int Dimensions, access::mode AccessMode,
           access::placeholder IsPlaceholder>
-class accessor<DataT, Dimensions, AccessMode, access::target::image_array,
-               IsPlaceholder>
+class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
+                                    access::target::image_array, IsPlaceholder>
     : public detail::image_accessor<DataT, Dimensions + 1, AccessMode,
                                     access::target::image, IsPlaceholder> {
 #ifdef __SYCL_DEVICE_ONLY__
diff --git a/sycl/include/CL/sycl/detail/defines.hpp b/sycl/include/CL/sycl/detail/defines.hpp
index e2461333209f7..ae5d9f2c7f3e9 100644
--- a/sycl/include/CL/sycl/detail/defines.hpp
+++ b/sycl/include/CL/sycl/detail/defines.hpp
@@ -20,3 +20,9 @@
 #warning "No assumptions will be emitted due to no __builtin_assume available"
 #endif
 #endif
+
+#if __has_attribute(sycl_special_class)
+#define __SYCL_SPECIAL_CLASS __attribute__((sycl_special_class))
+#else
+#define __SYCL_SPECIAL_CLASS
+#endif
diff --git a/sycl/include/CL/sycl/sampler.hpp b/sycl/include/CL/sycl/sampler.hpp
index 5ebae973bcd4e..256fa21be1cac 100644
--- a/sycl/include/CL/sycl/sampler.hpp
+++ b/sycl/include/CL/sycl/sampler.hpp
@@ -62,7 +62,7 @@ class sampler_impl;
 /// \sa sycl_api_acc
 ///
 /// \ingroup sycl_api
-class __SYCL_EXPORT sampler {
+class __SYCL_EXPORT __SYCL_SPECIAL_CLASS sampler {
 public:
   sampler(coordinate_normalization_mode normalizationMode,
           addressing_mode addressingMode, filtering_mode filteringMode,
diff --git a/sycl/include/CL/sycl/stream.hpp b/sycl/include/CL/sycl/stream.hpp
index 1bf9a83488c2f..c63c30d0d1f72 100644
--- a/sycl/include/CL/sycl/stream.hpp
+++ b/sycl/include/CL/sycl/stream.hpp
@@ -739,7 +739,7 @@ inline __width_manipulator__ setw(int Width) {
 /// vector and SYCL types to the console.
 ///
 /// \ingroup sycl_api
-class __SYCL_EXPORT stream {
+class __SYCL_EXPORT __SYCL_SPECIAL_CLASS stream {
 public:
 #ifdef __SYCL_DEVICE_ONLY__
   // Default constructor for objects later initialized with __init member.