diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index e0116feddc392..9d9e15553665c 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -359,6 +359,13 @@ class SYCLIntegrationHeader {
     Itr->updateKernelNames(Name, StableName);
   }
 
+  /// Signals that emission of __sycl_device_global_registration type and
+  /// declaration of variable __sycl_device_global_registrar of this type in
+  /// integration header is required.
+  void addDeviceGlobalRegistration() {
+    NeedToEmitDeviceGlobalRegistration = true;
+  }
+
 private:
   // Kernel actual parameter descriptor.
   struct KernelParamDesc {
@@ -433,6 +440,10 @@ class SYCLIntegrationHeader {
   llvm::SmallVector<SpecConstID, 4> SpecConsts;
 
   Sema &S;
+
+  /// Keeps track of whether declaration of __sycl_device_global_registration
+  /// type and __sycl_device_global_registrar variable are required to emit.
+  bool NeedToEmitDeviceGlobalRegistration = false;
 };
 
 class SYCLIntegrationFooter {
@@ -440,14 +451,12 @@ class SYCLIntegrationFooter {
   SYCLIntegrationFooter(Sema &S) : S(S) {}
   bool emit(StringRef MainSrc);
   void addVarDecl(const VarDecl *VD);
-  bool isDeviceGlobalsEmitted() { return DeviceGlobalsEmitted; }
 
 private:
   bool emit(raw_ostream &O);
   Sema &S;
   llvm::SmallVector<const VarDecl *> GlobalVars;
   void emitSpecIDName(raw_ostream &O, const VarDecl *VD);
-  bool DeviceGlobalsEmitted = false;
 };
 
 /// Tracks expected type during expression parsing, for use in code completion.
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 8174a24bbc27c..7ade4155c3285 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -4656,7 +4656,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   // whose sole purpose is to run its constructor before the application's
   // main() function.
 
-  if (S.getSyclIntegrationFooter().isDeviceGlobalsEmitted()) {
+  if (NeedToEmitDeviceGlobalRegistration) {
     O << "namespace {\n";
 
     O << "class __sycl_device_global_registration {\n";
@@ -5020,6 +5020,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
 
   llvm::SmallSet<const VarDecl *, 8> Visited;
   bool EmittedFirstSpecConstant = false;
+  bool DeviceGlobalsEmitted = false;
 
   // Used to uniquely name the 'shim's as we generate the names in each
   // anonymous namespace.
@@ -5103,6 +5104,8 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
     OS << "}\n";
     OS << "} // namespace (unnamed)\n";
     OS << "} // namespace sycl::detail\n";
+
+    S.getSyclIntegrationHeader().addDeviceGlobalRegistration();
   }
   return true;
 }
diff --git a/clang/test/CodeGenSYCL/int_header_without_kernels.cpp b/clang/test/CodeGenSYCL/int_header_without_kernels.cpp
new file mode 100644
index 0000000000000..14fad42a87ab5
--- /dev/null
+++ b/clang/test/CodeGenSYCL/int_header_without_kernels.cpp
@@ -0,0 +1,38 @@
+// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-int-footer=%t.footer.h -fsycl-int-header=%t.header.h %s -emit-llvm -o %t.ll
+// RUN: FileCheck -input-file=%t.footer.h %s --check-prefix=CHECK-FOOTER
+// RUN: FileCheck -input-file=%t.header.h %s --check-prefix=CHECK-HEADER
+
+// This test checks that integration header and footer are emitted correctly
+// for device_global variables even without kernels.
+
+#include "sycl.hpp"
+
+using namespace cl::sycl::ext::oneapi;
+
+// CHECK-HEADER: __SYCL_INLINE_NAMESPACE(cl) {
+// CHECK-HEADER-NEXT: namespace sycl {
+// CHECK-HEADER-NEXT: namespace detail {
+// CHECK-HEADER-NEXT: namespace {
+// CHECK-HEADER-NEXT: class __sycl_device_global_registration {
+// CHECK-HEADER-NEXT: public:
+// CHECK-HEADER-NEXT:   __sycl_device_global_registration() noexcept;
+// CHECK-HEADER-NEXT: };
+// CHECK-HEADER-NEXT: __sycl_device_global_registration __sycl_device_global_registrar;
+// CHECK-HEADER-NEXT: } // namespace
+// CHECK-HEADER: } // namespace detail
+// CHECK-HEADER: } // namespace sycl
+// CHECK-HEADER: } // __SYCL_INLINE_NAMESPACE(cl)
+
+// CHECK-FOOTER: #include <CL/sycl/detail/defines_elementary.hpp>
+
+// CHECK-FOOTER: #include <CL/sycl/detail/device_global_map.hpp>
+// CHECK-FOOTER: namespace sycl::detail {
+// CHECK-FOOTER-NEXT: namespace {
+// CHECK-FOOTER-NEXT: __sycl_device_global_registration::__sycl_device_global_registration() noexcept {
+
+device_global<int> Basic;
+// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::Basic, "_Z5Basic");
+
+// CHECK-FOOTER-NEXT: }
+// CHECK-FOOTER-NEXT: }
+// CHECK-FOOTER-NEXT: }