diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 05f33867af877..b544ff38a54d2 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1257,17 +1257,8 @@ def SYCLUsesAspects : InheritableAttr {
   let Subjects = SubjectList<[CXXRecord, Function], ErrorDiag>;
   let Args = [VariadicExprArgument<"Aspects">];
   let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
-  // Only used internally by the SYCL implementation
-  let Documentation = [Undocumented];
-}
-
-def SYCLDeviceGlobal : InheritableAttr {
-  let Spellings = [CXX11<"__sycl_detail__", "device_global">];
-  let Subjects = SubjectList<[CXXRecord], ErrorDiag>;
-  let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
-  // Only used internally by the SYCL implementation
+  // Only used internally by SYCL implementation
   let Documentation = [Undocumented];
-  let SimpleHandler = 1;
 }
 
 // Marks functions which must not be vectorized via horizontal SIMT widening,
@@ -1448,6 +1439,24 @@ def SYCLIntelMaxGlobalWorkDim : InheritableAttr {
   let SupportsNonconformingLambdaSyntax = 1;
 }
 
+def SYCLDeviceGlobal: InheritableAttr {
+  let Spellings = [CXX11<"__sycl_detail__", "device_global">];
+  let Subjects = SubjectList<[CXXRecord], ErrorDiag>;
+  let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
+  // Only used internally by SYCL implementation
+  let Documentation = [SYCLDeviceGlobalAttrDocs];
+  let SimpleHandler = 1;
+}
+
+def SYCLGlobalVariableAllowed : InheritableAttr {
+  let Spellings = [CXX11<"__sycl_detail__", "global_variable_allowed">];
+  let Subjects = SubjectList<[CXXRecord], ErrorDiag>;
+  let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
+  // Only used internally by SYCL implementation
+  let Documentation = [SYCLGlobalVariableAllowedAttrDocs];
+  let SimpleHandler = 1;
+}
+
 def SYCLIntelNoGlobalWorkOffset : InheritableAttr {
   let Spellings = [CXX11<"intel", "no_global_work_offset">];
   let Args = [ExprArgument<"Value", /*optional*/1>];
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 99d5660d36bea..462df1357fb58 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -3033,6 +3033,49 @@ function. In SYCL 2020 mode, the attribute is not propagated to the kernel.
   }];
 }
 
+def SYCLDeviceGlobalAttrDocs : Documentation {
+  let Category = DocCatType;
+  let Heading = "__sycl_detail__::device_global";
+  let Content = [{
+This attribute is part of support for SYCL device_global feature.
+[[__sycl_detail__::device_global]] attribute is used for checking restrictions
+on variable declarations using the device_global type instead of the class name.
+Global or static variables of type decorated with this attribute have
+`sycl-unique-id`, an LLVM IR attribute, added to the definition of each such
+variable, which provides a unique string identifier using
+__builtin_sycl_unique_stable_id.
+We do not intend to support this as a general attribute that user code can use,
+so we have this attribute in sycl_detail namespace.
+
+.. code-block:: c++
+
+  template<typename T>
+  struct [[__sycl_detail__::device_global]] device_global {}
+
+  device_global<int> Foo;
+  }];
+}
+
+def SYCLGlobalVariableAllowedAttrDocs : Documentation {
+  let Category = DocCatType;
+  let Heading = "__sycl_detail__::global_variable_allowed";
+  let Content = [{
+This attribute is part of support for SYCL device_global feature.
+[[__sycl_detail__::global_variable_allowed]] attribute is used to avoid
+diagnosing an error when global or static variables of type decorated with this
+attribute are referenced in device code. We do not intend to support this as a
+general attribute that user code can use, therefore it is wrapped in
+sycl_detail namespace.
+
+.. code-block:: c++
+  
+  template<typename T>
+  struct [[__sycl_detail__::global_variable_allowed]] device_global {}
+  
+  device_global<int> Foo;
+  }];
+}
+
 def SYCLFPGAPipeDocs : Documentation {
   let Category = DocCatStmt;
   let Heading = "pipe (read_only, write_only)";
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 18ef9b2fb7fa8..3c521df759a54 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -7094,6 +7094,11 @@ def warn_format_nonliteral : Warning<
   "format string is not a string literal">,
   InGroup<FormatNonLiteral>, DefaultIgnore;
 
+def err_sycl_device_global_incorrect_scope : Error<
+  "'device_global' variables must be static or declared at namespace scope">;
+def err_sycl_device_global_not_publicly_accessible: Error<
+  "'device_global' member variable %0 is not publicly accessible from namespace scope">;
+
 def err_unexpected_interface : Error<
   "unexpected interface name %0: expected expression">;
 def err_ref_non_value : Error<"%0 does not refer to a value">;
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 923d087dc01ac..e6119c24630e0 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -13163,6 +13163,19 @@ class Sema final {
                                    SourceLocation BuiltinLoc,
                                    SourceLocation RParenLoc);
 
+  template <typename AttrTy>
+  bool isTypeDecoratedWithDeclAttribute(QualType Ty) {
+    const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl();
+    if (!RecTy)
+      return false;
+    if (auto *CTSD = dyn_cast<ClassTemplateSpecializationDecl>(RecTy)) {
+      ClassTemplateDecl *Template = CTSD->getSpecializedTemplate();
+      if (CXXRecordDecl *RD = Template->getTemplatedDecl())
+        return RD->hasAttr<AttrTy>();
+    }
+    return RecTy->hasAttr<AttrTy>();
+  }
+
 private:
   bool SemaBuiltinPrefetch(CallExpr *TheCall);
   bool SemaBuiltinAllocaWithAlign(CallExpr *TheCall);
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index af1cbbc65b4d6..fe88f5c99f5c9 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -1594,10 +1594,17 @@ void CodeGenModule::EmitCtorList(CtorList &Fns, const char *GlobalName) {
   llvm::FunctionType* CtorFTy = llvm::FunctionType::get(VoidTy, false);
   llvm::Type *CtorPFTy = llvm::PointerType::get(CtorFTy,
       TheModule.getDataLayout().getProgramAddressSpace());
+  llvm::PointerType *TargetType = VoidPtrTy;
+  // Get target type when templated global variables are used,
+  // to emit them correctly in the target (default) address space and avoid
+  // emitting them in a private address space.
+  if (getLangOpts().SYCLIsDevice)
+    TargetType = llvm::IntegerType::getInt8PtrTy(
+        getLLVMContext(), getContext().getTargetAddressSpace(LangAS::Default));
 
   // Get the type of a ctor entry, { i32, void ()*, i8* }.
-  llvm::StructType *CtorStructTy = llvm::StructType::get(
-      Int32Ty, CtorPFTy, VoidPtrTy);
+  llvm::StructType *CtorStructTy =
+      llvm::StructType::get(Int32Ty, CtorPFTy, TargetType);
 
   // Construct the constructor and destructor arrays.
   ConstantInitBuilder builder(*this);
@@ -1606,10 +1613,12 @@ void CodeGenModule::EmitCtorList(CtorList &Fns, const char *GlobalName) {
     auto ctor = ctors.beginStruct(CtorStructTy);
     ctor.addInt(Int32Ty, I.Priority);
     ctor.add(llvm::ConstantExpr::getBitCast(I.Initializer, CtorPFTy));
+    // Emit appropriate bitcasts for pointers of different address spaces.
     if (I.AssociatedData)
-      ctor.add(llvm::ConstantExpr::getBitCast(I.AssociatedData, VoidPtrTy));
+      ctor.add(llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast(
+          I.AssociatedData, TargetType));
     else
-      ctor.addNullPointer(VoidPtrTy);
+      ctor.addNullPointer(TargetType);
     ctor.finishAndAddTo(ctors);
   }
 
@@ -2428,19 +2437,26 @@ static void emitUsed(CodeGenModule &CGM, StringRef Name,
   // Don't create llvm.used if there is no need.
   if (List.empty())
     return;
+  // For SYCL emit pointers in the default address space which is a superset of
+  // other address spaces, so that casts from any other address spaces will be
+  // valid.
+  llvm::PointerType *TargetType = CGM.Int8PtrTy;
+  if (CGM.getLangOpts().SYCLIsDevice)
+    TargetType = llvm::IntegerType::getInt8PtrTy(
+        CGM.getLLVMContext(),
+        CGM.getContext().getTargetAddressSpace(LangAS::Default));
 
   // Convert List to what ConstantArray needs.
   SmallVector<llvm::Constant*, 8> UsedArray;
   UsedArray.resize(List.size());
   for (unsigned i = 0, e = List.size(); i != e; ++i) {
-    UsedArray[i] =
-        llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast(
-            cast<llvm::Constant>(&*List[i]), CGM.Int8PtrTy);
+    UsedArray[i] = llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast(
+        cast<llvm::Constant>(&*List[i]), TargetType);
   }
 
   if (UsedArray.empty())
     return;
-  llvm::ArrayType *ATy = llvm::ArrayType::get(CGM.Int8PtrTy, UsedArray.size());
+  llvm::ArrayType *ATy = llvm::ArrayType::get(TargetType, UsedArray.size());
 
   auto *GV = new llvm::GlobalVariable(
       CGM.getModule(), ATy, false, llvm::GlobalValue::AppendingLinkage,
@@ -2846,6 +2862,15 @@ void CodeGenModule::AddGlobalAnnotations(const ValueDecl *D,
     Annotations.push_back(EmitAnnotateAttr(GV, I, D->getLocation()));
 }
 
+// Add "sycl-unique-id" llvm IR attribute that has a unique string generated
+// by __builtin_sycl_unique_stable_id for global variables marked with
+// SYCL device_global attribute.
+static void addSYCLUniqueID(llvm::GlobalVariable *GV, const VarDecl *VD,
+                            ASTContext &Context) {
+  auto builtinString = SYCLUniqueStableIdExpr::ComputeName(Context, VD);
+  GV->addAttribute("sycl-unique-id", builtinString);
+}
+
 bool CodeGenModule::isInNoSanitizeList(SanitizerMask Kind, llvm::Function *Fn,
                                        SourceLocation Loc) const {
   const auto &NoSanitizeL = getContext().getNoSanitizeList();
@@ -4942,6 +4967,14 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
   if (getLangOpts().SYCLIsDevice)
     addGlobalIntelFPGAAnnotation(D, GV);
 
+  // If VarDecl has a type decorated with SYCL device_global attribute, emit IR
+  // attribute 'sycl-unique-id'.
+  if (getLangOpts().SYCLIsDevice) {
+    const RecordDecl *RD = D->getType()->getAsRecordDecl();
+    if (RD && RD->hasAttr<SYCLDeviceGlobalAttr>())
+      addSYCLUniqueID(GV, D, Context);
+  }
+
   if (D->getType().isRestrictQualified()) {
     llvm::LLVMContext &Context = getLLVMContext();
 
diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
index 9a20b20a14942..1c5df431e07dd 100644
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -1703,7 +1703,9 @@ class DeferredDiagnosticsEmitter
   void visitUsedDecl(SourceLocation Loc, Decl *D) {
     if (S.LangOpts.SYCLIsDevice && ShouldEmitRootNode) {
       if (auto *VD = dyn_cast<VarDecl>(D)) {
-        if (!S.checkAllowedSYCLInitializer(VD)) {
+        if (!S.checkAllowedSYCLInitializer(VD) &&
+            !S.isTypeDecoratedWithDeclAttribute<SYCLGlobalVariableAllowedAttr>(
+                VD->getType())) {
           S.Diag(Loc, diag::err_sycl_restrict)
               << Sema::KernelConstStaticVariable;
           return;
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 27d7d0d5c5900..cb0eb43d8c8b6 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -7400,12 +7400,23 @@ NamedDecl *Sema::ActOnVariableDeclarator(
       NewVD->setTSCSpec(TSCS);
   }
 
-  // Static variables declared inside SYCL device code must be const or
-  // constexpr
-  if (getLangOpts().SYCLIsDevice)
-    if (SCSpec == DeclSpec::SCS_static && !R.isConstant(Context))
+  // Global variables with types decorated with device_global attribute must be
+  // static if they are declared in SYCL device code.
+  if (getLangOpts().SYCLIsDevice) {
+    if (SCSpec != DeclSpec::SCS_static && !NewVD->hasGlobalStorage() &&
+        isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
+            NewVD->getType()))
+      Diag(D.getIdentifierLoc(), diag::err_sycl_device_global_incorrect_scope);
+
+    // Static variables declared inside SYCL device code must be const or
+    // constexpr unless their types are decorated with global_variable_allowed
+    // attribute.
+    if (SCSpec == DeclSpec::SCS_static && !R.isConstant(Context) &&
+        !isTypeDecoratedWithDeclAttribute<SYCLGlobalVariableAllowedAttr>(
+            NewVD->getType()))
       SYCLDiagIfDeviceCode(D.getIdentifierLoc(), diag::err_sycl_restrict)
           << Sema::KernelNonConstStaticDataVariable;
+  }
 
   switch (D.getDeclSpec().getConstexprSpecifier()) {
   case ConstexprSpecKind::Unspecified:
diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp
index 90a1cf514eb6e..a1e763ee73cbb 100644
--- a/clang/lib/Sema/SemaDeclCXX.cpp
+++ b/clang/lib/Sema/SemaDeclCXX.cpp
@@ -3529,6 +3529,19 @@ Sema::ActOnCXXMemberDeclarator(Scope *S, AccessSpecifier AS, Declarator &D,
     }
   }
 
+  // Emit diagnostic if a private member of type decorated with device_global
+  // attribute is accessed.
+  if (getLangOpts().SYCLIsDevice) {
+    if (auto Value = dyn_cast<ValueDecl>(Member)) {
+      if (isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
+              Value->getType()) &&
+          Value->getAccess() != AS_public) {
+        Diag(Loc, diag::err_sycl_device_global_not_publicly_accessible)
+            << Value;
+      }
+    }
+  }
+
   return Member;
 }
 
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index 85ccd66d247b8..eecef13454109 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -227,16 +227,20 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
           (!isUnevaluatedContext() && !isConstantEvaluated());
       bool IsEsimdPrivateGlobal = isSYCLEsimdPrivateGlobal(VD);
       // Non-const statics are not allowed in SYCL except for ESIMD or with the
-      // SYCLGlobalVar attribute.
+      // SYCLGlobalVar or SYCLGlobalVariableAllowed attribute.
       if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst &&
           VD->getStorageClass() == SC_Static &&
-          !VD->hasAttr<SYCLGlobalVarAttr>())
+          !VD->hasAttr<SYCLGlobalVarAttr>() &&
+          !isTypeDecoratedWithDeclAttribute<SYCLGlobalVariableAllowedAttr>(
+              VD->getType()))
         SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict)
             << Sema::KernelNonConstStaticDataVariable;
       // Non-const globals are not allowed in SYCL except for ESIMD or with the
-      // SYCLGlobalVar attribute.
+      // SYCLGlobalVar or SYCLGlobalVariableAllowed attribute.
       else if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst &&
-               VD->hasGlobalStorage() && !VD->hasAttr<SYCLGlobalVarAttr>())
+               VD->hasGlobalStorage() && !VD->hasAttr<SYCLGlobalVarAttr>() &&
+               !isTypeDecoratedWithDeclAttribute<SYCLGlobalVariableAllowedAttr>(
+                   VD->getType()))
         SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict)
             << Sema::KernelGlobalVariable;
       // ESIMD globals cannot be used in a SYCL context.
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 4a94c25bbfd17..0db660f1dfd8b 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -123,10 +123,6 @@ class Util {
   /// specialization id class.
   static bool isSyclSpecIdType(QualType Ty);
 
-  /// Checks whether given clang type is a full specialization of the SYCL
-  /// device_global class.
-  static bool isSyclDeviceGlobalType(QualType Ty);
-
   /// Checks whether given clang type is a full specialization of the SYCL
   /// kernel_handler class.
   static bool isSyclKernelHandlerType(QualType Ty);
@@ -4896,7 +4892,8 @@ void SYCLIntegrationFooter::addVarDecl(const VarDecl *VD) {
     return;
   // Step 1: ensure that this is of the correct type template specialization.
   if (!Util::isSyclSpecIdType(VD->getType()) &&
-      !Util::isSyclDeviceGlobalType(VD->getType())) {
+      !S.isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
+          VD->getType())) {
     // Handle the case where this could be a deduced type, such as a deduction
     // guide. We have to do this here since this function, unlike most of the
     // rest of this file, is called during Sema instead of after it. We will
@@ -5076,7 +5073,8 @@ 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()) &&
-        !Util::isSyclDeviceGlobalType(VD->getType()))
+        !S.isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
+            VD->getType()))
       continue;
 
     // Skip if we've already visited this.
@@ -5090,7 +5088,8 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
 
     Visited.insert(VD);
     std::string TopShim = EmitShims(OS, ShimCounter, Policy, VD);
-    if (Util::isSyclDeviceGlobalType(VD->getType())) {
+    if (S.isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
+            VD->getType())) {
       DeviceGlobalsEmitted = true;
       DeviceGlobOS << "device_global_map::add(";
       DeviceGlobOS << "(void *)&";
@@ -5189,18 +5188,6 @@ bool Util::isSyclSpecIdType(QualType Ty) {
   return matchQualifiedTypeName(Ty, Scopes);
 }
 
-bool Util::isSyclDeviceGlobalType(QualType Ty) {
-  const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl();
-  if (!RecTy)
-    return false;
-  if (auto *CTSD = dyn_cast<ClassTemplateSpecializationDecl>(RecTy)) {
-    ClassTemplateDecl *Template = CTSD->getSpecializedTemplate();
-    if (CXXRecordDecl *RD = Template->getTemplatedDecl())
-      return RD->hasAttr<SYCLDeviceGlobalAttr>();
-  }
-  return RecTy->hasAttr<SYCLDeviceGlobalAttr>();
-}
-
 bool Util::isSyclKernelHandlerType(QualType Ty) {
   std::array<DeclContextDesc, 3> Scopes = {
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
index facf6dfa2fd16..7200b51695d2d 100644
--- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -130,11 +130,13 @@ struct no_alias {
 };
 } // namespace property
 
+// device_global type decorated with attributes
 template <typename T>
-class [[__sycl_detail__::device_global]] device_global {
+class [[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] device_global {
 public:
   const T &get() const noexcept { return *Data; }
   device_global() {}
+  operator T &() noexcept { return *Data; }
 
 private:
   T *Data;
diff --git a/clang/test/CodeGenSYCL/device_global.cpp b/clang/test/CodeGenSYCL/device_global.cpp
new file mode 100644
index 0000000000000..c1b68db290282
--- /dev/null
+++ b/clang/test/CodeGenSYCL/device_global.cpp
@@ -0,0 +1,103 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -fsycl-unique-prefix=THE_PREFIX -std=c++17 -emit-llvm %s -o - | FileCheck %s
+#include "sycl.hpp"
+
+// Test cases below show that 'sycl-unique-id' LLVM IR attribute is attached to the
+// global variable whose type is decorated with device_global attribute, and that a
+// unique string is generated.
+
+using namespace sycl::ext::oneapi;
+using namespace cl::sycl;
+queue q;
+
+device_global<int> A;
+static device_global<int> B;
+
+struct Foo {
+  static device_global<int> C;
+};
+device_global<int> Foo::C;
+// CHECK: @A = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[A_ATTRS:[0-9]+]]
+// CHECK: @_ZL1B = internal addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[B_ATTRS:[0-9]+]]
+// CHECK: @_ZN3Foo1CE = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[C_ATTRS:[0-9]+]]
+
+device_global<int> same_name;
+namespace NS {
+device_global<int> same_name;
+}
+// CHECK: @same_name = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[SAME_NAME_ATTRS:[0-9]+]]
+// CHECK: @_ZN2NS9same_nameE = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[SAME_NAME_NS_ATTRS:[0-9]+]]
+
+// decorated with only global_variable_allowed attribute
+template <typename T>
+class [[__sycl_detail__::global_variable_allowed]] only_global_var_allowed {
+public:
+  const T &get() const noexcept { return *Data; }
+  only_global_var_allowed() {}
+  operator T &() noexcept { return *Data; }
+
+private:
+  T *Data;
+};
+
+// check that we don't generate `sycl-unique-id` IR attribute if class does not use
+// [[__sycl_detail__::device_global]]
+only_global_var_allowed<int> no_device_global;
+// CHECK: @no_device_global = addrspace(1) global %class.only_global_var_allowed zeroinitializer, align 8{{$}}
+
+inline namespace Bar {
+device_global<float> InlineNS;
+}
+// CHECK: @_ZN3Bar8InlineNSE = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8 #[[BAR_INLINENS_ATTRS:[0-9]+]]
+
+template <typename T> struct TS {
+public:
+  static device_global<T> d;
+};
+template <> device_global<int> TS<int>::d{};
+// CHECK: @_ZN2TSIiE1dE = addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[TEMPLATED_WRAPPER_ATTRS:[0-9]+]]
+
+template <typename T>
+device_global<T> templ_dev_global;
+// CHECK: @[[TEMPL_DEV_GLOB:[a-zA-Z0-9_]+]] = linkonce_odr addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, comdat, align 8 #[[TEMPL_DEV_GLOB_ATTRS:[0-9]+]]
+
+void foo() {
+  q.submit([&](handler &h) {
+    h.single_task<class kernel_name_1>([=]() {
+      (void)A;
+      (void)B;
+      (void)Foo::C;
+      (void)same_name;
+      (void)NS::same_name;
+      (void)no_device_global;
+      (void)Bar::InlineNS;
+      auto AA = TS<int>::d.get();
+      auto val = templ_dev_global<int>.get();
+    });
+  });
+}
+
+namespace {
+device_global<int> same_name;
+}
+// CHECK: @_ZN12_GLOBAL__N_19same_nameE = internal addrspace(1) global %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 8 #[[SAME_NAME_ANON_NS_ATTRS:[0-9]+]]
+
+namespace {
+void bar() {
+  q.submit([&](handler &h) {
+    h.single_task<class kernel_name>([=]() { int A = same_name; });
+  });
+}
+} // namespace
+
+// CHECK: @llvm.global_ctors = appending global [2 x { i32, void ()*, i8 addrspace(4)* }] [{ i32, void ()*, i8 addrspace(4)* } { i32 65535, void ()* @__cxx_global_var_init{{.*}}, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @[[TEMPL_DEV_GLOB]] to i8 addrspace(1)*) to i8 addrspace(4)*) }, { i32, void ()*, i8 addrspace(4)* } { i32 65535, void ()* @_GLOBAL__sub_I_device_global.cpp, i8 addrspace(4)* null }]
+// CHECK: @llvm.used = appending global [1 x i8 addrspace(4)*] [i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @[[TEMPL_DEV_GLOB]] to i8 addrspace(1)*) to i8 addrspace(4)*)], section "llvm.metadata"
+
+// CHECK: attributes #[[A_ATTRS]] = { "sycl-unique-id"="_Z1A" }
+// CHECK: attributes #[[B_ATTRS]] = { "sycl-unique-id"="THE_PREFIX____ZL1B" }
+// CHECK: attributes #[[C_ATTRS]] = { "sycl-unique-id"="_ZN3Foo1CE" }
+// CHECK: attributes #[[SAME_NAME_ATTRS]] = { "sycl-unique-id"="_Z9same_name" }
+// CHECK: attributes #[[SAME_NAME_NS_ATTRS]] = { "sycl-unique-id"="_ZN2NS9same_nameE" }
+// CHECK: attributes #[[BAR_INLINENS_ATTRS]] = { "sycl-unique-id"="_ZN3Bar8InlineNSE" }
+// CHECK: attributes #[[TEMPLATED_WRAPPER_ATTRS]] = { "sycl-unique-id"="_ZN2TSIiE1dE" }
+// CHECK: attributes #[[TEMPL_DEV_GLOB_ATTRS]] = { "sycl-unique-id"="_Z16templ_dev_globalIiE" }
+// CHECK: attributes #[[SAME_NAME_ANON_NS_ATTRS]] = { "sycl-unique-id"="THE_PREFIX____ZN12_GLOBAL__N_19same_nameE" }
diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
index 4c1c3e14bca59..59908bbe226f7 100644
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -160,6 +160,7 @@
 // CHECK-NEXT: SYCLDeviceGlobal (SubjectMatchRule_record)
 // CHECK-NEXT: SYCLDeviceHas (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLDeviceIndirectlyCallable (SubjectMatchRule_function)
+// CHECK-NEXT: SYCLGlobalVariableAllowed (SubjectMatchRule_record)
 // CHECK-NEXT: SYCLIntelFPGADisableLoopPipelining (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLIntelFPGAInitiationInterval (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLIntelFPGAMaxConcurrency (SubjectMatchRule_function)
diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp
index 67c3f1aeb8ce0..abc4358d739ff 100644
--- a/clang/test/SemaSYCL/Inputs/sycl.hpp
+++ b/clang/test/SemaSYCL/Inputs/sycl.hpp
@@ -65,6 +65,18 @@ namespace ext {
 namespace oneapi {
 template <typename... properties>
 class accessor_property_list {};
+
+// device_global type decorated with attributes
+template <typename T>
+struct [[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] device_global {
+public:
+  const T &get() const noexcept { return *Data; }
+  device_global() {}
+  operator T &() noexcept { return *Data; }
+
+private:
+  T *Data;
+};
 } // namespace oneapi
 } // namespace ext
 
@@ -353,7 +365,7 @@ template <typename T, access::address_space AS> class multi_ptr {
   pointer_t m_Pointer;
 
 public:
-  multi_ptr(T *Ptr) : m_Pointer((pointer_t)(Ptr)) {}
+  multi_ptr(T *Ptr) : m_Pointer((pointer_t)(Ptr)) {} // #MultiPtrConstructor
   pointer_t get() { return m_Pointer; }
 };
 
diff --git a/clang/test/SemaSYCL/device_global.cpp b/clang/test/SemaSYCL/device_global.cpp
new file mode 100644
index 0000000000000..66296eaf923dd
--- /dev/null
+++ b/clang/test/SemaSYCL/device_global.cpp
@@ -0,0 +1,63 @@
+// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -verify %s
+#include "Inputs/sycl.hpp"
+
+// Test cases below check for valid usage of device_global and
+// global_variable_allowed attributes, and that they are being correctly
+// generated in the AST.
+using namespace sycl::ext::oneapi;
+
+device_global<int> glob;                  // OK
+static device_global<float> static_glob;  // OK
+inline device_global<double> inline_glob; // OK
+static const device_global<int> static_const_glob;
+
+struct Foo {
+  static device_global<char> d; // OK
+};
+device_global<char> Foo::d;
+
+struct Baz {
+private:
+  // expected-error@+1{{'device_global' member variable 'f' is not publicly accessible from namespace scope}}
+  static device_global<int> f;
+};
+device_global<int> Baz::f;
+
+device_global<int[4]> not_array; // OK
+
+device_global<int> same_name; // OK
+namespace foo {
+device_global<int> same_name; // OK
+}
+namespace {
+device_global<int> same_name; // OK
+}
+
+// expected-error@+2{{'device_global' attribute only applies to classes}}
+// expected-error@+1{{'global_variable_allowed' attribute only applies to classes}}
+[[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] int integer;
+
+// expected-error@+2{{'device_global' attribute only applies to classes}}
+// expected-error@+1{{'global_variable_allowed' attribute only applies to classes}}
+[[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] int *pointer;
+
+union [[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] a_union;
+
+int main() {
+  cl::sycl::kernel_single_task<class KernelName1>([=]() {
+    (void)glob;
+    (void)static_glob;
+    (void)inline_glob;
+    (void)static_const_glob;
+    (void)Foo::d;
+  });
+
+  cl::sycl::kernel_single_task<class KernelName2>([]() {
+    // expected-error@+1{{'device_global' variables must be static or declared at namespace scope}}
+    device_global<int> non_static;
+
+    // expect no error on non_const_static declaration if decorated with
+    // [[__sycl_detail__::global_variable_allowed]]
+    static device_global<int> non_const_static;
+  });
+}
diff --git a/clang/test/SemaSYCL/device_global_ast.cpp b/clang/test/SemaSYCL/device_global_ast.cpp
new file mode 100644
index 0000000000000..e39a41ab7fa24
--- /dev/null
+++ b/clang/test/SemaSYCL/device_global_ast.cpp
@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -ast-dump %s | FileCheck %s
+#include "Inputs/sycl.hpp"
+
+// Test cases below check that DeviceGlobalAttr and GlobalVariableAllowedAttr
+// are correctly emitted.
+
+using namespace sycl::ext::oneapi;
+
+device_global<int> glob;
+// CHECK: ClassTemplateDecl {{.*}} device_global
+// CHECK: CXXRecordDecl {{.*}} struct device_global definition
+// CHECK: SYCLDeviceGlobalAttr {{.*}}
+// CHECK: SYCLGlobalVariableAllowedAttr {{.*}}
+// CHECK: ClassTemplateSpecializationDecl {{.*}} struct device_global definition
+// CHECK: SYCLDeviceGlobalAttr {{.*}}
+// CHECK: SYCLGlobalVariableAllowedAttr {{.*}}
diff --git a/clang/test/SemaSYCL/explicit-cast-to-generic.cpp b/clang/test/SemaSYCL/explicit-cast-to-generic.cpp
index 6fdcf14f3e296..f8a1583a8669b 100644
--- a/clang/test/SemaSYCL/explicit-cast-to-generic.cpp
+++ b/clang/test/SemaSYCL/explicit-cast-to-generic.cpp
@@ -15,7 +15,7 @@ void __attribute__((sycl_device)) onDeviceUsages(multi_ptr<int, access::address_
   // expected-warning@+1 {{explicit cast from 'int *' to '__local int *' potentially leads to an invalid address space cast in the resulting code}}
   __attribute__((opencl_local)) int *LL = (__attribute((opencl_local)) int *)NoAS;
 
-  // expected-warning@Inputs/sycl.hpp:356 {{explicit cast from 'int *' to 'sycl::multi_ptr<int, sycl::access::address_space::private_space>::pointer_t' (aka '__private int *') potentially leads to an invalid address space cast in the resulting code}}
+  // expected-warning@#MultiPtrConstructor {{explicit cast from 'int *' to 'sycl::multi_ptr<int, sycl::access::address_space::private_space>::pointer_t' (aka '__private int *') potentially leads to an invalid address space cast in the resulting code}}
   // expected-note@+1 {{called by 'onDeviceUsages'}}
   auto P = multi_ptr<int, access::address_space::private_space>{F.get()};