diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 6bc17f6eceea0..7e7b41dde9391 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1281,6 +1281,23 @@ def SYCLSpecialClass: InheritableAttr {
   let Documentation = [SYCLSpecialClassDocs];
 }
 
+def SYCLType: InheritableAttr {
+  let Spellings = [CXX11<"__sycl_detail__", "sycl_type">];
+  let Subjects = SubjectList<[CXXRecord, Enum], ErrorDiag>;
+  let LangOpts = [SYCLIsDevice, SYCLIsHost];
+  let Args = [EnumArgument<"Type", "SYCLType",
+                           ["accessor", "local_accessor", "spec_constant",
+                            "specialization_id", "kernel_handler", "buffer_location",
+                            "no_alias", "accessor_property_list", "group",
+                            "private_memory", "aspect"],
+                           ["accessor", "local_accessor", "spec_constant",
+                            "specialization_id", "kernel_handler", "buffer_location",
+                            "no_alias", "accessor_property_list", "group",
+                            "private_memory", "aspect"]>];
+  // Only used internally by SYCL implementation
+  let Documentation = [InternalOnly];
+}
+
 def SYCLDeviceHas : InheritableAttr {
   let Spellings = [CXX11<"sycl", "device_has">];
   let Subjects = SubjectList<[Function], ErrorDiag>;
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index d87176ec939c8..7c397c9b7245f 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -4098,6 +4098,8 @@ def warn_transparent_union_attribute_zero_fields : Warning<
 def warn_attribute_type_not_supported : Warning<
   "%0 attribute argument not supported: %1">,
   InGroup<IgnoredAttributes>;
+def err_attribute_argument_not_supported : Error<
+  "%0 attribute argument %1 is not supported">;
 def warn_attribute_unknown_visibility : Warning<"unknown visibility %0">,
   InGroup<IgnoredAttributes>;
 def warn_attribute_protected_visibility :
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index a61ad4da77950..165b22dab5a0f 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -10858,6 +10858,9 @@ class Sema final {
   ReqdWorkGroupSizeAttr *
   MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A);
 
+  SYCLTypeAttr *MergeSYCLTypeAttr(Decl *D, const AttributeCommonInfo &CI,
+                                  SYCLTypeAttr::SYCLType TypeName);
+
   /// Only called on function definitions; if there is a MSVC #pragma optimize
   /// in scope, consider changing the function's attributes based on the
   /// optimization list passed to the pragma.
@@ -13546,12 +13549,16 @@ class Sema final {
     const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl();
     if (!RecTy)
       return false;
+
+    if (RecTy->hasAttr<AttrTy>())
+      return true;
+
     if (auto *CTSD = dyn_cast<ClassTemplateSpecializationDecl>(RecTy)) {
       ClassTemplateDecl *Template = CTSD->getSpecializedTemplate();
       if (CXXRecordDecl *RD = Template->getTemplatedDecl())
         return RD->hasAttr<AttrTy>();
     }
-    return RecTy->hasAttr<AttrTy>();
+    return false;
   }
 
 private:
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index f337d0f2df173..e38295a2f6d04 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -2918,6 +2918,8 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D,
     NewAttr = S.MergeSYCLDeviceHasAttr(D, *A);
   else if (const auto *A = dyn_cast<SYCLUsesAspectsAttr>(Attr))
     NewAttr = S.MergeSYCLUsesAspectsAttr(D, *A);
+  else if (const auto *A = dyn_cast<SYCLTypeAttr>(Attr))
+    NewAttr = S.MergeSYCLTypeAttr(D, *A, A->getType());
   else if (const auto *A = dyn_cast<SYCLIntelPipeIOAttr>(Attr))
     NewAttr = S.MergeSYCLIntelPipeIOAttr(D, *A);
   else if (const auto *A = dyn_cast<SYCLIntelMaxWorkGroupSizeAttr>(Attr))
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 959ba38d80930..7097d4223182c 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -10426,46 +10426,15 @@ static void handleFunctionReturnThunksAttr(Sema &S, Decl *D,
   D->addAttr(FunctionReturnThunksAttr::Create(S.Context, Kind, AL));
 }
 
-static constexpr std::pair<Decl::Kind, StringRef>
-MakeDeclContextDesc(Decl::Kind K, StringRef SR) {
-  return std::pair<Decl::Kind, StringRef>{K, SR};
-}
-
-// FIXME: Refactor Util class in SemaSYCL.cpp to avoid following
-// code duplication.
 bool isDeviceAspectType(const QualType Ty) {
   const EnumType *ET = Ty->getAs<EnumType>();
   if (!ET)
     return false;
 
-  std::array<std::pair<Decl::Kind, StringRef>, 3> Scopes = {
-      MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
-      MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
-      MakeDeclContextDesc(Decl::Kind::Enum, "aspect")};
-
-  const auto *Ctx = cast<DeclContext>(ET->getDecl());
-  StringRef Name = "";
-
-  for (const auto &Scope : llvm::reverse(Scopes)) {
-    Decl::Kind DK = Ctx->getDeclKind();
-    if (DK != Scope.first)
-      return false;
+  if (const auto *Attr = ET->getDecl()->getAttr<SYCLTypeAttr>())
+    return Attr->getType() == SYCLTypeAttr::aspect;
 
-    switch (DK) {
-    case Decl::Kind::Enum:
-      Name = cast<EnumDecl>(Ctx)->getName();
-      break;
-    case Decl::Kind::Namespace:
-      Name = cast<NamespaceDecl>(Ctx)->getName();
-      break;
-    default:
-      llvm_unreachable("isDeviceAspectType: decl kind not supported");
-    }
-    if (Name != Scope.second)
-      return false;
-    Ctx = Ctx->getParent();
-  }
-  return Ctx->isTranslationUnit();
+  return false;
 }
 
 SYCLDeviceHasAttr *Sema::MergeSYCLDeviceHasAttr(Decl *D,
@@ -10590,6 +10559,39 @@ static void handleSYCLKernelAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   handleSimpleAttribute<SYCLKernelAttr>(S, D, AL);
 }
 
+SYCLTypeAttr *Sema::MergeSYCLTypeAttr(Decl *D, const AttributeCommonInfo &CI,
+                                      SYCLTypeAttr::SYCLType TypeName) {
+  if (const auto *ExistingAttr = D->getAttr<SYCLTypeAttr>()) {
+    if (ExistingAttr->getType() != TypeName) {
+      Diag(ExistingAttr->getLoc(), diag::err_duplicate_attribute)
+          << ExistingAttr;
+      Diag(CI.getLoc(), diag::note_previous_attribute);
+    }
+    // Do not add duplicate attribute
+    return nullptr;
+  }
+  return ::new (Context) SYCLTypeAttr(Context, CI, TypeName);
+}
+
+static void handleSYCLTypeAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
+  if (!AL.isArgIdent(0)) {
+    S.Diag(AL.getLoc(), diag::err_attribute_argument_type)
+        << AL << AANT_ArgumentIdentifier;
+    return;
+  }
+
+  IdentifierInfo *II = AL.getArgAsIdent(0)->Ident;
+  SYCLTypeAttr::SYCLType Type;
+
+  if (!SYCLTypeAttr::ConvertStrToSYCLType(II->getName(), Type)) {
+    S.Diag(AL.getLoc(), diag::err_attribute_argument_not_supported) << AL << II;
+    return;
+  }
+
+  if (SYCLTypeAttr *NewAttr = S.MergeSYCLTypeAttr(D, AL, Type))
+    D->addAttr(NewAttr);
+}
+
 static void handleDestroyAttr(Sema &S, Decl *D, const ParsedAttr &A) {
   if (!cast<VarDecl>(D)->hasGlobalStorage()) {
     S.Diag(D->getLocation(), diag::err_destroy_attr_on_non_static_var)
@@ -11142,6 +11144,9 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL,
   case ParsedAttr::AT_SYCLSpecialClass:
     handleSimpleAttribute<SYCLSpecialClassAttr>(S, D, AL);
     break;
+  case ParsedAttr::AT_SYCLType:
+    handleSYCLTypeAttr(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 c1424fa65705e..a763109d39aad 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -67,84 +67,45 @@ static constexpr llvm::StringLiteral LibstdcxxFailedAssertion =
     "__failed_assertion";
 constexpr unsigned MaxKernelArgsSize = 2048;
 
-namespace {
+static bool isSyclType(QualType Ty, SYCLTypeAttr::SYCLType TypeName) {
+  const auto *RD = Ty->getAsCXXRecordDecl();
+  if (!RD)
+    return false;
 
-/// Various utilities.
-class Util {
-public:
-  using DeclContextDesc = std::pair<Decl::Kind, StringRef>;
-
-  template <size_t N>
-  static constexpr DeclContextDesc MakeDeclContextDesc(Decl::Kind K,
-                                                       const char (&Str)[N]) {
-    // FIXME: This SHOULD be able to use the StringLiteral constructor here
-    // instead, however this seems to fail with an 'invalid string literal' note
-    // on the correct constructor in some build configurations.  We need to
-    // figure that out before reverting this to use the StringLiteral
-    // constructor.
-    return DeclContextDesc{K, StringRef{Str, N - 1}};
-  }
-
-  static constexpr DeclContextDesc MakeDeclContextDesc(Decl::Kind K,
-                                                       StringRef SR) {
-    return DeclContextDesc{K, SR};
-  }
-
-  static bool isSyclSpecialType(QualType Ty);
-
-  /// Checks whether given clang type is a full specialization of the SYCL
-  /// accessor_property_list class.
-  static bool isAccessorPropertyListType(QualType Ty);
-
-  /// Checks whether given clang type is a full specialization of the SYCL
-  /// no_alias class.
-  static bool isSyclAccessorNoAliasPropertyType(QualType Ty);
-
-  /// Checks whether given clang type is a full specialization of the SYCL
-  /// buffer_location class.
-  static bool isSyclBufferLocationType(QualType Ty);
-
-  /// Checks whether given clang type is a standard SYCL API class with given
-  /// name.
-  /// \param Ty    the clang type being checked
-  /// \param Name  the class name checked against
-  /// \param Tmpl  whether the class is template instantiation or simple record
-  static bool isSyclType(QualType Ty, StringRef Name, bool Tmpl = false);
-
-  /// Checks whether given clang type is a standard SYCL API accessor class,
-  /// the check assumes the type is templated.
-  /// \param Ty    the clang type being checked
-  static bool isSyclAccessorType(QualType Ty);
-
-  /// Checks whether given clang type is a full specialization of the SYCL
-  /// specialization constant class.
-  static bool isSyclSpecConstantType(QualType Ty);
-
-  /// Checks whether given clang type is a full specialization of the SYCL
-  /// specialization id class.
-  static bool isSyclSpecIdType(QualType Ty);
-
-  /// Checks whether given clang type is a full specialization of the SYCL
-  /// kernel_handler class.
-  static bool isSyclKernelHandlerType(QualType Ty);
-
-  // Checks declaration context hierarchy.
-  /// \param DC     the context of the item to be checked.
-  /// \param Scopes the declaration scopes leading from the item context to the
-  ///               translation unit (excluding the latter)
-  static bool matchContext(const DeclContext *DC,
-                           ArrayRef<Util::DeclContextDesc> Scopes);
-
-  /// Checks whether given clang type is declared in the given hierarchy of
-  /// declaration contexts.
-  /// \param Ty         the clang type being checked
-  /// \param Scopes     the declaration scopes leading from the type to the
-  ///     translation unit (excluding the latter)
-  static bool matchQualifiedTypeName(QualType Ty,
-                                     ArrayRef<Util::DeclContextDesc> Scopes);
-};
+  if (const auto *Attr = RD->getAttr<SYCLTypeAttr>())
+    return Attr->getType() == TypeName;
+
+  if (const auto *CTSD = dyn_cast<ClassTemplateSpecializationDecl>(RD))
+    if (CXXRecordDecl *TemplateDecl =
+            CTSD->getSpecializedTemplate()->getTemplatedDecl())
+      if (const auto *Attr = TemplateDecl->getAttr<SYCLTypeAttr>())
+        return Attr->getType() == TypeName;
+
+  return false;
+}
+
+static bool isSyclAccessorType(QualType Ty) {
+  return isSyclType(Ty, SYCLTypeAttr::accessor) ||
+         isSyclType(Ty, SYCLTypeAttr::local_accessor);
+}
+
+// FIXME: Accessor property lists should be modified to use compile-time
+// properties. Once implemented, this function (and possibly all/most code
+// in SemaSYCL.cpp handling no_alias and buffer_location property) can be
+// removed.
+static bool isAccessorPropertyType(QualType Ty,
+                                   SYCLTypeAttr::SYCLType TypeName) {
+  if (const auto *RD = Ty->getAsCXXRecordDecl())
+    if (const auto *Parent = dyn_cast<CXXRecordDecl>(RD->getParent()))
+      if (const auto *Attr = Parent->getAttr<SYCLTypeAttr>())
+        return Attr->getType() == TypeName;
+
+  return false;
+}
 
-} // anonymous namespace
+static bool isSyclSpecialType(QualType Ty, Sema &S) {
+  return S.isTypeDecoratedWithDeclAttribute<SYCLSpecialClassAttr>(Ty);
+}
 
 ExprResult Sema::ActOnSYCLBuiltinNumFieldsExpr(ParsedType PT) {
   TypeSourceInfo *TInfo = nullptr;
@@ -929,7 +890,7 @@ class MarkWIScopeFnVisitor : public RecursiveASTVisitor<MarkWIScopeFnVisitor> {
       // not a direct call - continue search
       return true;
     QualType Ty = Ctx.getRecordType(Call->getRecordDecl());
-    if (!Util::isSyclType(Ty, "group", true /*Tmpl*/))
+    if (!isSyclType(Ty, SYCLTypeAttr::group))
       // not a member of sycl::group - continue search
       return true;
     auto Name = Callee->getName();
@@ -949,7 +910,7 @@ class MarkWIScopeFnVisitor : public RecursiveASTVisitor<MarkWIScopeFnVisitor> {
 };
 
 static bool isSYCLPrivateMemoryVar(VarDecl *VD) {
-  return Util::isSyclType(VD->getType(), "private_memory", true /*Tmpl*/);
+  return isSyclType(VD->getType(), SYCLTypeAttr::private_memory);
 }
 
 static void addScopeAttrToLocalVars(CXXMethodDecl &F) {
@@ -1028,7 +989,7 @@ static ParamDesc makeParamDesc(ASTContext &Ctx, StringRef Name, QualType Ty) {
 /// \return the target of given SYCL accessor type
 static target getAccessTarget(QualType FieldTy,
                               const ClassTemplateSpecializationDecl *AccTy) {
-  if (Util::isSyclType(FieldTy, "local_accessor", true /*Tmpl*/))
+  if (isSyclType(FieldTy, SYCLTypeAttr::local_accessor))
     return local;
 
   return static_cast<target>(
@@ -1073,7 +1034,7 @@ static ParmVarDecl *getSyclKernelHandlerArg(FunctionDecl *KernelCallerFunc) {
   // Specialization constants in SYCL 2020 are not captured by lambda and
   // accessed through new optional lambda argument kernel_handler
   auto IsHandlerLambda = [](ParmVarDecl *PVD) {
-    return Util::isSyclKernelHandlerType(PVD->getType());
+    return isSyclType(PVD->getType(), SYCLTypeAttr::kernel_handler);
   };
 
   assert(llvm::count_if(KernelCallerFunc->parameters(), IsHandlerLambda) <= 1 &&
@@ -1195,10 +1156,10 @@ class KernelObjVisitor {
     for (const auto &Base : Range) {
       QualType BaseTy = Base.getType();
       // Handle accessor class as base
-      if (Util::isSyclSpecialType(BaseTy)) {
+      if (isSyclSpecialType(BaseTy, SemaRef))
         (void)std::initializer_list<int>{
             (Handlers.handleSyclSpecialType(Owner, Base, BaseTy), 0)...};
-      } else
+      else
         // For all other bases, visit the record
         visitRecord(Owner, Base, BaseTy->getAsCXXRecordDecl(), BaseTy,
                     Handlers...);
@@ -1277,9 +1238,9 @@ class KernelObjVisitor {
   template <typename... HandlerTys>
   void visitField(const CXXRecordDecl *Owner, FieldDecl *Field,
                   QualType FieldTy, HandlerTys &... Handlers) {
-    if (Util::isSyclSpecialType(FieldTy))
+    if (isSyclSpecialType(FieldTy, SemaRef))
       KF_FOR_EACH(handleSyclSpecialType, Field, FieldTy);
-    else if (Util::isSyclSpecConstantType(FieldTy))
+    else if (isSyclType(FieldTy, SYCLTypeAttr::spec_constant))
       KF_FOR_EACH(handleSyclSpecConstantType, Field, FieldTy);
     else if (FieldTy->isStructureOrClassType()) {
       if (KF_FOR_EACH(handleStructType, Field, FieldTy)) {
@@ -1562,7 +1523,7 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
           Loc, diag::err_sycl_invalid_accessor_property_template_param);
 
     QualType PropListTy = PropList.getAsType();
-    if (!Util::isAccessorPropertyListType(PropListTy))
+    if (!isSyclType(PropListTy, SYCLTypeAttr::accessor_property_list))
       return SemaRef.Diag(
           Loc, diag::err_sycl_invalid_accessor_property_template_param);
 
@@ -1588,7 +1549,7 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
                    diag::err_sycl_invalid_accessor_property_list_template_param)
                << /*accessor_property_list pack argument*/ 1 << /*type*/ 1;
       QualType PropTy = Prop->getAsType();
-      if (Util::isSyclBufferLocationType(PropTy) &&
+      if (isAccessorPropertyType(PropTy, SYCLTypeAttr::buffer_location) &&
           checkBufferLocationType(PropTy, Loc))
         return true;
     }
@@ -1621,10 +1582,10 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
   }
 
   bool checkSyclSpecialType(QualType Ty, SourceRange Loc) {
-    assert(Util::isSyclSpecialType(Ty) &&
+    assert(isSyclSpecialType(Ty, SemaRef) &&
            "Should only be called on sycl special class types.");
     const RecordDecl *RecD = Ty->getAsRecordDecl();
-    if (IsSIMD && !Util::isSyclAccessorType(Ty))
+    if (IsSIMD && !isSyclAccessorType(Ty))
       return SemaRef.Diag(Loc.getBegin(),
                           diag::err_sycl_esimd_not_supported_for_type)
              << RecD;
@@ -1902,9 +1863,9 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     for (TemplateArgument::pack_iterator Prop = TemplArg.pack_begin();
          Prop != TemplArg.pack_end(); ++Prop) {
       QualType PropTy = Prop->getAsType();
-      if (Util::isSyclAccessorNoAliasPropertyType(PropTy))
+      if (isAccessorPropertyType(PropTy, SYCLTypeAttr::no_alias))
         handleNoAliasProperty(Param, PropTy, Loc);
-      if (Util::isSyclBufferLocationType(PropTy))
+      if (isAccessorPropertyType(PropTy, SYCLTypeAttr::buffer_location))
         handleBufferLocationProperty(Param, PropTy, Loc);
     }
   }
@@ -1941,7 +1902,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     handleAccessorPropertyList(Params.back(), RecordDecl, Loc);
 
     // If "accessor" type check if read only
-    if (Util::isSyclType(FieldTy, "accessor", true /*Tmpl*/)) {
+    if (isSyclType(FieldTy, SYCLTypeAttr::accessor)) {
       // Get access mode of accessor.
       const auto *AccessorSpecializationDecl =
           cast<ClassTemplateSpecializationDecl>(RecordDecl);
@@ -1966,7 +1927,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     const auto *RecordDecl = FieldTy->getAsCXXRecordDecl();
     assert(RecordDecl && "The type must be a RecordDecl");
     llvm::StringLiteral MethodName =
-        KernelDecl->hasAttr<SYCLSimdAttr>() && Util::isSyclAccessorType(FieldTy)
+        KernelDecl->hasAttr<SYCLSimdAttr>() && isSyclAccessorType(FieldTy)
             ? InitESIMDMethodName
             : InitMethodName;
     CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName);
@@ -1990,8 +1951,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
       // 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::isSyclAccessorType(FieldTy))
+      if (ParamTy.getTypePtr()->isPointerType() && isSyclAccessorType(FieldTy))
         handleAccessorType(FieldTy, RecordDecl, FD->getBeginLoc());
     }
     LastParamIndex = ParamIndex;
@@ -2086,7 +2046,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     const auto *RecordDecl = FieldTy->getAsCXXRecordDecl();
     assert(RecordDecl && "The type must be a RecordDecl");
     llvm::StringLiteral MethodName =
-        KernelDecl->hasAttr<SYCLSimdAttr>() && Util::isSyclAccessorType(FieldTy)
+        KernelDecl->hasAttr<SYCLSimdAttr>() && isSyclAccessorType(FieldTy)
             ? InitESIMDMethodName
             : InitMethodName;
     CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName);
@@ -2104,8 +2064,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
       // 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::isSyclAccessorType(FieldTy))
+      if (ParamTy.getTypePtr()->isPointerType() && isSyclAccessorType(FieldTy))
         handleAccessorType(FieldTy, RecordDecl, BS.getBeginLoc());
     }
     LastParamIndex = ParamIndex;
@@ -2226,9 +2185,9 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler {
   bool handleSpecialType(QualType FieldTy) {
     const CXXRecordDecl *RecordDecl = FieldTy->getAsCXXRecordDecl();
     assert(RecordDecl && "The type must be a RecordDecl");
-    llvm::StringLiteral MethodName =
-        (IsSIMD && Util::isSyclAccessorType(FieldTy)) ? InitESIMDMethodName
-                                                      : InitMethodName;
+    llvm::StringLiteral MethodName = (IsSIMD && isSyclAccessorType(FieldTy))
+                                         ? InitESIMDMethodName
+                                         : InitMethodName;
     CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName);
     assert(InitMethod && "The type must have the __init method");
     for (const ParmVarDecl *Param : InitMethod->parameters())
@@ -3145,7 +3104,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
   bool handleSyclSpecialType(FieldDecl *FD, QualType FieldTy) final {
     const auto *ClassTy = FieldTy->getAsCXXRecordDecl();
     assert(ClassTy && "Type must be a C++ record type");
-    if (Util::isSyclAccessorType(FieldTy)) {
+    if (isSyclAccessorType(FieldTy)) {
       const auto *AccTy =
           cast<ClassTemplateSpecializationDecl>(FieldTy->getAsRecordDecl());
       assert(AccTy->getTemplateArgs().size() >= 2 &&
@@ -4872,7 +4831,7 @@ void SYCLIntegrationFooter::addVarDecl(const VarDecl *VD) {
   if (isa<VarTemplatePartialSpecializationDecl>(VD))
     return;
   // Step 1: ensure that this is of the correct type template specialization.
-  if (!Util::isSyclSpecIdType(VD->getType()) &&
+  if (!isSyclType(VD->getType(), SYCLTypeAttr::specialization_id) &&
       !S.isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
           VD->getType())) {
     // Handle the case where this could be a deduced type, such as a deduction
@@ -5054,7 +5013,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
 
     // Skip if this isn't a SpecIdType or DeviceGlobal.  This can happen if it
     // was a deduced type.
-    if (!Util::isSyclSpecIdType(VD->getType()) &&
+    if (!isSyclType(VD->getType(), SYCLTypeAttr::specialization_id) &&
         !S.isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
             VD->getType()))
       continue;
@@ -5130,137 +5089,3 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
   }
   return true;
 }
-
-// -----------------------------------------------------------------------------
-// Utility class methods
-// -----------------------------------------------------------------------------
-bool Util::isSyclSpecialType(const QualType Ty) {
-  const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl();
-  if (!RecTy)
-    return false;
-  return RecTy->hasAttr<SYCLSpecialClassAttr>();
-}
-
-bool Util::isSyclSpecConstantType(QualType Ty) {
-  std::array<DeclContextDesc, 6> Scopes = {
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ext"),
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "oneapi"),
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "experimental"),
-      Util::MakeDeclContextDesc(Decl::Kind::ClassTemplateSpecialization,
-                                "spec_constant")};
-  return matchQualifiedTypeName(Ty, Scopes);
-}
-
-bool Util::isSyclSpecIdType(QualType Ty) {
-  std::array<DeclContextDesc, 3> Scopes = {
-      Util::MakeDeclContextDesc(clang::Decl::Kind::Namespace, "sycl"),
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
-      Util::MakeDeclContextDesc(Decl::Kind::ClassTemplateSpecialization,
-                                "specialization_id")};
-  return matchQualifiedTypeName(Ty, Scopes);
-}
-
-bool Util::isSyclKernelHandlerType(QualType Ty) {
-  std::array<DeclContextDesc, 3> Scopes = {
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
-      Util::MakeDeclContextDesc(Decl::Kind::CXXRecord, "kernel_handler")};
-  return matchQualifiedTypeName(Ty, Scopes);
-}
-
-bool Util::isSyclAccessorNoAliasPropertyType(QualType Ty) {
-  std::array<DeclContextDesc, 7> Scopes = {
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ext"),
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "oneapi"),
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "property"),
-      Util::MakeDeclContextDesc(Decl::Kind::CXXRecord, "no_alias"),
-      Util::MakeDeclContextDesc(Decl::Kind::ClassTemplateSpecialization,
-                                "instance")};
-  return matchQualifiedTypeName(Ty, Scopes);
-}
-
-bool Util::isSyclBufferLocationType(QualType Ty) {
-  std::array<DeclContextDesc, 7> Scopes = {
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ext"),
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "intel"),
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "property"),
-      Util::MakeDeclContextDesc(Decl::Kind::CXXRecord, "buffer_location"),
-      Util::MakeDeclContextDesc(Decl::Kind::ClassTemplateSpecialization,
-                                "instance")};
-  return matchQualifiedTypeName(Ty, Scopes);
-}
-
-bool Util::isSyclType(QualType Ty, StringRef Name, bool Tmpl) {
-  Decl::Kind ClassDeclKind =
-      Tmpl ? Decl::Kind::ClassTemplateSpecialization : Decl::Kind::CXXRecord;
-  std::array<DeclContextDesc, 3> Scopes = {
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
-      Util::MakeDeclContextDesc(ClassDeclKind, Name)};
-  return matchQualifiedTypeName(Ty, Scopes);
-}
-
-bool Util::isSyclAccessorType(QualType Ty) {
-  return isSyclType(Ty, "accessor", true /* Tmpl */) ||
-         isSyclType(Ty, "local_accessor", true /* Tmpl */);
-}
-
-bool Util::isAccessorPropertyListType(QualType Ty) {
-  std::array<DeclContextDesc, 5> Scopes = {
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ext"),
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "oneapi"),
-      Util::MakeDeclContextDesc(Decl::Kind::ClassTemplateSpecialization,
-                                "accessor_property_list")};
-  return matchQualifiedTypeName(Ty, Scopes);
-}
-
-bool Util::matchContext(const DeclContext *Ctx,
-                        ArrayRef<Util::DeclContextDesc> Scopes) {
-  // The idea: check the declaration context chain starting from the item
-  // itself. At each step check the context is of expected kind
-  // (namespace) and name.
-  StringRef Name = "";
-
-  for (const auto &Scope : llvm::reverse(Scopes)) {
-    Decl::Kind DK = Ctx->getDeclKind();
-    if (DK != Scope.first)
-      return false;
-
-    switch (DK) {
-    case Decl::Kind::ClassTemplateSpecialization:
-      // ClassTemplateSpecializationDecl inherits from CXXRecordDecl
-    case Decl::Kind::CXXRecord:
-      Name = cast<CXXRecordDecl>(Ctx)->getName();
-      break;
-    case Decl::Kind::Namespace:
-      Name = cast<NamespaceDecl>(Ctx)->getName();
-      break;
-    default:
-      llvm_unreachable("matchContext: decl kind not supported");
-    }
-    if (Name != Scope.second)
-      return false;
-    Ctx = Ctx->getParent();
-  }
-  return Ctx->isTranslationUnit() ||
-         (Ctx->isExternCXXContext() &&
-          Ctx->getEnclosingNamespaceContext()->isTranslationUnit());
-}
-
-bool Util::matchQualifiedTypeName(QualType Ty,
-                                  ArrayRef<Util::DeclContextDesc> Scopes) {
-  const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl();
-
-  if (!RecTy)
-    return false; // only classes/structs supported
-  const auto *Ctx = cast<DeclContext>(RecTy);
-  return Util::matchContext(Ctx, Scopes);
-}
diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
index a7424df625128..b6f06b1500304 100644
--- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -1,6 +1,7 @@
 #pragma once
 
 #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
+#define __SYCL_TYPE(x) [[__sycl_detail__::sycl_type(x)]]
 
 extern "C" int printf(const char* fmt, ...);
 
@@ -29,7 +30,7 @@ class __attribute__((sycl_special_class)) sampler {
 };
 
 template <int dimensions = 1>
-class group {
+class __SYCL_TYPE(group) group {
 public:
   group() = default; // fake constructor
 };
@@ -69,7 +70,7 @@ enum class address_space : int {
 } // namespace access
 
 // Dummy aspect enum with limited enumerators
-enum class aspect {
+enum class __SYCL_TYPE(aspect) aspect {
   host = 0,
   cpu = 1,
   gpu = 2,
@@ -119,7 +120,7 @@ namespace ext {
 namespace intel {
 namespace property {
 // Compile time known accessor property
-struct buffer_location {
+struct __SYCL_TYPE(buffer_location) buffer_location {
   template <int> class instance {};
 };
 } // namespace property
@@ -130,7 +131,7 @@ namespace ext {
 namespace oneapi {
 namespace property {
 // Compile time known accessor property
-struct no_alias {
+struct __SYCL_TYPE(no_alias) no_alias {
   template <bool> class instance {};
 };
 } // namespace property
@@ -172,7 +173,7 @@ template <int dim> struct item {
 namespace ext {
 namespace oneapi {
 template <typename... properties>
-class accessor_property_list {};
+class __SYCL_TYPE(accessor_property_list) accessor_property_list {};
 } // namespace oneapi
 } // namespace ext
 
@@ -201,7 +202,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 __attribute__((sycl_special_class)) accessor {
+class __attribute__((sycl_special_class)) __SYCL_TYPE(accessor) accessor {
 
 public:
   void use(void) const {}
@@ -267,7 +268,7 @@ struct _ImageImplT {
 };
 
 template <typename dataT, int dimensions, access::mode accessmode>
-class __attribute__((sycl_special_class)) accessor<dataT, dimensions, accessmode, access::target::image, access::placeholder::false_t> {
+class __attribute__((sycl_special_class)) __SYCL_TYPE(accessor) accessor<dataT, dimensions, accessmode, access::target::image, access::placeholder::false_t> {
 public:
   void use(void) const {}
   template <typename... T>
@@ -292,7 +293,7 @@ class accessor<dataT, dimensions, accessmode, access::target::host_image, access
 };
 
 template <typename dataT, int dimensions>
-class __attribute__((sycl_special_class))
+class __attribute__((sycl_special_class)) __SYCL_TYPE(local_accessor)
 local_accessor: public accessor<dataT,
         dimensions, access::mode::read_write,
         access::target::local> {
@@ -344,7 +345,7 @@ namespace ext {
 namespace oneapi {
 namespace experimental {
 template <typename T, typename ID = T>
-class spec_constant {
+class __SYCL_TYPE(spec_constant) spec_constant {
 public:
   spec_constant() {}
   spec_constant(T Cst) {}
@@ -375,11 +376,11 @@ int printf(const __SYCL_CONSTANT_AS char *__format, Args... args) {
 } // namespace oneapi
 } // namespace ext
 
-class kernel_handler {
+class __SYCL_TYPE(kernel_handler) kernel_handler {
   void __init_specialization_constants_buffer(char *specialization_constants_buffer) {}
 };
 
-template <typename T> class specialization_id {
+template <typename T> class __SYCL_TYPE(specialization_id) specialization_id {
 public:
   using value_type = T;
 
diff --git a/clang/test/CodeGenSYCL/accessor-readonly-invalid-lib.cpp b/clang/test/CodeGenSYCL/accessor-readonly-invalid-lib.cpp
index 2ca60450d5c56..7ca8bbb86a205 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 __attribute__((sycl_special_class)) accessor {
+class __attribute__((sycl_special_class)) [[__sycl_detail__::sycl_type(accessor)]] 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 113a178ece83f..f249468f88cf4 100644
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -183,6 +183,7 @@
 // CHECK-NEXT: SYCLRegisterNum (SubjectMatchRule_variable_is_global)
 // CHECK-NEXT: SYCLSimd (SubjectMatchRule_function, SubjectMatchRule_variable_is_global)
 // CHECK-NEXT: SYCLSpecialClass (SubjectMatchRule_record)
+// CHECK-NEXT: SYCLType (SubjectMatchRule_record, SubjectMatchRule_enum)
 // CHECK-NEXT: SYCLUsesAspects (SubjectMatchRule_record, SubjectMatchRule_function)
 // CHECK-NEXT: ScopedLockable (SubjectMatchRule_record)
 // CHECK-NEXT: Section (SubjectMatchRule_function, SubjectMatchRule_variable_is_global, SubjectMatchRule_objc_method, SubjectMatchRule_objc_property)
diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp
index 19adf15143e85..dd2c08d774c62 100644
--- a/clang/test/SemaSYCL/Inputs/sycl.hpp
+++ b/clang/test/SemaSYCL/Inputs/sycl.hpp
@@ -1,6 +1,8 @@
 #ifndef SYCL_HPP
 #define SYCL_HPP
 
+#define __SYCL_TYPE(x) [[__sycl_detail__::sycl_type(x)]]
+
 // Shared code for SYCL tests
 
 namespace sycl {
@@ -39,7 +41,7 @@ enum class address_space : int {
 } // namespace access
 
 // Dummy aspect enum with limited enumerators
-enum class aspect {
+enum class __SYCL_TYPE(aspect) aspect {
   host = 0,
   cpu = 1,
   gpu = 2,
@@ -54,7 +56,7 @@ class property_list {};
 namespace ext {
 namespace intel {
 namespace property {
-struct buffer_location {
+struct __SYCL_TYPE(buffer_location) buffer_location {
   template <int> class instance {};
 };
 } // namespace property
@@ -64,7 +66,7 @@ struct buffer_location {
 namespace ext {
 namespace oneapi {
 template <typename... properties>
-class accessor_property_list {};
+class __SYCL_TYPE(accessor_property_list) accessor_property_list {};
 
 // device_global type decorated with attributes
 template <typename T>
@@ -131,7 +133,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 __attribute__((sycl_special_class)) accessor {
+class __attribute__((sycl_special_class)) __SYCL_TYPE(accessor) accessor {
 public:
   void use(void) const {}
   void use(void *) const {}
@@ -194,7 +196,7 @@ struct _ImageImplT {
 };
 
 template <typename dataT, int dimensions, access::mode accessmode>
-class __attribute__((sycl_special_class)) accessor<dataT, dimensions, accessmode, access::target::image, access::placeholder::false_t> {
+class __attribute__((sycl_special_class))  __SYCL_TYPE(accessor) accessor<dataT, dimensions, accessmode, access::target::image, access::placeholder::false_t> {
 public:
   void use(void) const {}
   template <typename... T>
@@ -208,7 +210,7 @@ class __attribute__((sycl_special_class)) accessor<dataT, dimensions, accessmode
 };
 
 template <typename dataT, int dimensions>
-class __attribute__((sycl_special_class))
+class __attribute__((sycl_special_class)) __SYCL_TYPE(local_accessor)
 local_accessor: public accessor<dataT,
         dimensions, access::mode::read_write,
         access::target::local> {
@@ -260,12 +262,12 @@ struct get_kernel_name_t<auto_name, Type> {
 };
 
 template <int dimensions = 1>
-class group {
+class __SYCL_TYPE(group) group {
 public:
   group() = default; // fake constructor
 };
 
-class kernel_handler {
+class __SYCL_TYPE(kernel_handler) kernel_handler {
   void __init_specialization_constants_buffer(char *specialization_constants_buffer) {}
 };
 
@@ -393,7 +395,7 @@ namespace ext {
 namespace oneapi {
 namespace experimental {
 template <typename T, typename ID = T>
-class spec_constant {
+class __SYCL_TYPE(spec_constant) spec_constant {
 public:
   spec_constant() {}
   explicit constexpr spec_constant(T defaultVal) : DefaultValue(defaultVal) {}
diff --git a/clang/test/SemaSYCL/sycl-type-attr-ast.cpp b/clang/test/SemaSYCL/sycl-type-attr-ast.cpp
new file mode 100644
index 0000000000000..71de3a67667da
--- /dev/null
+++ b/clang/test/SemaSYCL/sycl-type-attr-ast.cpp
@@ -0,0 +1,33 @@
+// RUN: %clang_cc1 -fsycl-is-device -ast-dump %s | FileCheck %s
+
+// Tests for AST of sycl_type() attribute
+
+class [[__sycl_detail__::sycl_type(accessor)]] accessor {};
+// CHECK: CXXRecordDecl {{.*}} class accessor definition
+// CHECK: SYCLTypeAttr {{.*}} accessor
+
+class accessor;
+// CHECK: CXXRecordDecl {{.*}} prev {{.*}} class accessor
+// CHECK: SYCLTypeAttr {{.*}} Inherited accessor
+
+enum class [[__sycl_detail__::sycl_type(aspect)]] aspect {};
+// CHECK: EnumDecl {{.*}} class aspect 'int'
+// CHECK: SYCLTypeAttr {{.*}} aspect
+
+template <typename T>
+class [[__sycl_detail__::sycl_type(local_accessor)]] local_accessor {};
+// CHECK: ClassTemplateDecl {{.*}} local_accessor
+// CHECK: CXXRecordDecl {{.*}} class local_accessor definition
+// CHECK: SYCLTypeAttr {{.*}} local_accessor
+
+class [[__sycl_detail__::sycl_type(spec_constant)]]
+[[__sycl_detail__::sycl_type(spec_constant)]] spec_constant;
+// CHECK: CXXRecordDecl {{.*}} spec_constant
+// CHECK: SYCLTypeAttr {{.*}} spec_constant
+// CHECK-NOT: SYCLTypeAttr {{.*}} spec_constant
+
+template <>
+class [[__sycl_detail__::sycl_type(local_accessor)]] local_accessor <int> {};
+// CHECK: ClassTemplateSpecializationDecl {{.*}} class local_accessor definition
+// CHECK: SYCLTypeAttr {{.*}} local_accessor
+
diff --git a/clang/test/SemaSYCL/sycl-type-attr.cpp b/clang/test/SemaSYCL/sycl-type-attr.cpp
new file mode 100644
index 0000000000000..fc41ab48cb44f
--- /dev/null
+++ b/clang/test/SemaSYCL/sycl-type-attr.cpp
@@ -0,0 +1,46 @@
+// RUN: %clang_cc1 -fsycl-is-device -verify %s
+
+// Diagnostic tests for sycl_type() attribute
+
+// expected-error@+1{{'sycl_type' attribute only applies to classes}}
+[[__sycl_detail__::sycl_type(accessor)]] int a;
+
+// expected-error@+1{{'sycl_type' attribute only applies to classes}}
+[[__sycl_detail__::sycl_type(accessor)]] void func1() {}
+
+// expected-error@+1{{'sycl_type' attribute requires an identifier}}
+class [[__sycl_detail__::sycl_type("accessor")]] A {};
+
+// expected-error@+1{{'sycl_type' attribute takes one argument}}
+class [[__sycl_detail__::sycl_type()]] B {};
+
+// expected-error@+1{{'sycl_type' attribute argument 'NotValidType' is not supported}}
+class [[__sycl_detail__::sycl_type(NotValidType)]] C {};
+
+// expected-note@+1{{previous attribute is here}}
+class [[__sycl_detail__::sycl_type(spec_constant)]] spec_constant;
+// expected-error@+1{{attribute 'sycl_type' is already applied with different arguments}}
+class [[__sycl_detail__::sycl_type(accessor)]] spec_constant {};
+
+// expected-error@+2{{attribute 'sycl_type' is already applied with different arguments}}
+// expected-note@+1{{previous attribute is here}}
+class [[__sycl_detail__::sycl_type(group)]] [[__sycl_detail__::sycl_type(accessor)]] group {};
+
+// Valid usage -
+
+class [[__sycl_detail__::sycl_type(accessor)]] accessor {};
+
+template <typename T>
+class [[__sycl_detail__::sycl_type(local_accessor)]] local_accessor {};
+
+enum class [[__sycl_detail__::sycl_type(aspect)]] aspect {};
+
+// No diagnostic for matching arguments. 
+class [[__sycl_detail__::sycl_type(kernel_handler)]] kernel_handler;
+class [[__sycl_detail__::sycl_type(kernel_handler)]]
+[[__sycl_detail__::sycl_type(kernel_handler)]] kernel_handler {};
+class kernel_handler;
+
+
+
+
diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index 6d9fd6a36ad9c..2ae84070fd3a1 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/sycl/include/sycl/accessor.hpp
@@ -789,7 +789,7 @@ class __image_array_slice__ {
 template <typename DataT, int Dimensions, access::mode AccessMode,
           access::target AccessTarget, access::placeholder IsPlaceholder,
           typename PropertyListT>
-class __SYCL_SPECIAL_CLASS accessor :
+class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
 #ifndef __SYCL_DEVICE_ONLY__
     public detail::AccessorBaseHost,
 #endif
@@ -2313,7 +2313,7 @@ class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
 };
 
 template <typename DataT, int Dimensions = 1>
-class __SYCL_SPECIAL_CLASS local_accessor
+class __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor
     : public local_accessor_base<DataT, Dimensions, access::mode::read_write,
                                  access::placeholder::false_t> {
 
@@ -2352,8 +2352,9 @@ class __SYCL_SPECIAL_CLASS local_accessor
 /// \ingroup sycl_api_acc
 template <typename DataT, int Dimensions, access::mode AccessMode,
           access::placeholder IsPlaceholder>
-class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
-                                    access::target::image, IsPlaceholder>
+class __SYCL_SPECIAL_CLASS
+__SYCL_TYPE(accessor) accessor<DataT, Dimensions, AccessMode,
+                               access::target::image, IsPlaceholder>
     : public detail::image_accessor<DataT, Dimensions, AccessMode,
                                     access::target::image, IsPlaceholder> {
 public:
@@ -2440,8 +2441,9 @@ 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 __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
-                                    access::target::image_array, IsPlaceholder>
+class __SYCL_SPECIAL_CLASS
+__SYCL_TYPE(accessor) 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/sycl/aspects.hpp b/sycl/include/sycl/aspects.hpp
index d71a903ac916b..ed7957f96e14b 100644
--- a/sycl/include/sycl/aspects.hpp
+++ b/sycl/include/sycl/aspects.hpp
@@ -12,7 +12,7 @@
 namespace sycl {
 __SYCL_INLINE_VER_NAMESPACE(_V1) {
 
-enum class aspect {
+enum class __SYCL_TYPE(aspect) aspect {
   host = 0,
   cpu = 1,
   gpu = 2,
diff --git a/sycl/include/sycl/detail/defines.hpp b/sycl/include/sycl/detail/defines.hpp
index 40cf18c953178..23df19ce4b284 100644
--- a/sycl/include/sycl/detail/defines.hpp
+++ b/sycl/include/sycl/detail/defines.hpp
@@ -26,3 +26,9 @@
 #else
 #define __SYCL_SPECIAL_CLASS
 #endif
+
+#if __has_cpp_attribute(__sycl_detail__::sycl_type)
+#define __SYCL_TYPE(x) [[__sycl_detail__::sycl_type(x)]]
+#else
+#define __SYCL_TYPE(x)
+#endif
diff --git a/sycl/include/sycl/ext/oneapi/accessor_property_list.hpp b/sycl/include/sycl/ext/oneapi/accessor_property_list.hpp
index bd1bbeb9b0592..4b0e384f7a378 100644
--- a/sycl/include/sycl/ext/oneapi/accessor_property_list.hpp
+++ b/sycl/include/sycl/ext/oneapi/accessor_property_list.hpp
@@ -41,7 +41,8 @@ template <typename T> struct is_compile_time_property : std::false_type {};
 ///
 /// \ingroup sycl_api
 template <typename... PropsT>
-class accessor_property_list : protected sycl::detail::PropertyListBase {
+class __SYCL_TYPE(accessor_property_list) accessor_property_list
+    : protected sycl::detail::PropertyListBase {
   // These structures check if compile-time-constant property is present in
   // list. For runtime properties this check is always true.
   template <class T, class U> struct AreSameTemplate : std::is_same<T, U> {};
diff --git a/sycl/include/sycl/ext/oneapi/experimental/spec_constant.hpp b/sycl/include/sycl/ext/oneapi/experimental/spec_constant.hpp
index e587cd4cb5bff..dd7a4ff9becd2 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/spec_constant.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/spec_constant.hpp
@@ -34,7 +34,8 @@ class spec_const_error : public compile_program_error {
 template <typename T, typename ID = T>
 class __SYCL2020_DEPRECATED(
     "Specialization constats extension is deprecated, use SYCL 2020"
-    " specialization constants instead") spec_constant {
+    " specialization constants instead")
+    __SYCL_TYPE(spec_constant) spec_constant {
 public:
   spec_constant() {}
 
diff --git a/sycl/include/sycl/group.hpp b/sycl/include/sycl/group.hpp
index 0f65326419bd2..fb2e0b5f491a7 100644
--- a/sycl/include/sycl/group.hpp
+++ b/sycl/include/sycl/group.hpp
@@ -49,7 +49,8 @@ static inline void workGroupBarrier() {
 //   class can be used to wrap the data. This class very simply constructs
 //   private data for a given group across the entire group.The id of the
 //   current work-item is passed to any access to grab the correct data.
-template <typename T, int Dimensions = 1> class private_memory {
+template <typename T, int Dimensions = 1>
+class __SYCL_TYPE(private_memory) private_memory {
 public:
   // Construct based directly off the number of work-items
   private_memory(const group<Dimensions> &G) {
@@ -90,7 +91,7 @@ template <typename T, int Dimensions = 1> class private_memory {
 /// within a parallel execution.
 ///
 /// \ingroup sycl_api
-template <int Dimensions = 1> class group {
+template <int Dimensions = 1> class __SYCL_TYPE(group) group {
 public:
 #ifndef __DISABLE_SYCL_INTEL_GROUP_ALGORITHMS__
   using id_type = id<Dimensions>;
diff --git a/sycl/include/sycl/kernel_handler.hpp b/sycl/include/sycl/kernel_handler.hpp
index f80fd6df8405d..e37888e44d34c 100644
--- a/sycl/include/sycl/kernel_handler.hpp
+++ b/sycl/include/sycl/kernel_handler.hpp
@@ -19,7 +19,7 @@ __SYCL_INLINE_VER_NAMESPACE(_V1) {
 /// Reading the value of a specialization constant
 ///
 /// \ingroup sycl_api
-class kernel_handler {
+class __SYCL_TYPE(kernel_handler) kernel_handler {
 public:
 #if __cplusplus >= 201703L
   template <auto &S>
diff --git a/sycl/include/sycl/properties/accessor_properties.hpp b/sycl/include/sycl/properties/accessor_properties.hpp
index cb721d4d1f079..5b5832bb51f02 100644
--- a/sycl/include/sycl/properties/accessor_properties.hpp
+++ b/sycl/include/sycl/properties/accessor_properties.hpp
@@ -48,7 +48,7 @@ constexpr const auto &noinit __SYCL2020_DEPRECATED("spelling is now: no_init") =
 namespace ext {
 namespace intel {
 namespace property {
-struct buffer_location {
+struct __SYCL_TYPE(buffer_location) buffer_location {
   template <int A = 0> struct instance {
     template <int B>
     constexpr bool operator==(const buffer_location::instance<B> &) const {
@@ -82,7 +82,7 @@ struct no_offset {
     }
   };
 };
-struct no_alias {
+struct __SYCL_TYPE(no_alias) no_alias {
   template <bool B = true> struct instance {
     constexpr bool operator==(const no_alias::instance<B> &) const {
       return true;
diff --git a/sycl/include/sycl/specialization_id.hpp b/sycl/include/sycl/specialization_id.hpp
index 8a94d324f6a6c..8570becba9e49 100644
--- a/sycl/include/sycl/specialization_id.hpp
+++ b/sycl/include/sycl/specialization_id.hpp
@@ -14,7 +14,7 @@ __SYCL_INLINE_VER_NAMESPACE(_V1) {
 /// Declaring a specialization constant
 ///
 /// \ingroup sycl_api
-template <typename T> class specialization_id {
+template <typename T> class __SYCL_TYPE(specialization_id) specialization_id {
 public:
   using value_type = T;