diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 996aca877350c..3b0c5bc9affe6 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -7302,22 +7302,30 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   }
 
   if (FreeFunctionCount > 0) {
+    // GlobalMapUpdater has to be in an anonymous namespace.
+    // Otherwise, if multiple translation units include the same integration
+    // header, there will be multiple varying definitions of GlobalMapUpdater
+    // with the same name across translation units, violating the C++'s One
+    // Definition Rule. Putting it in an anonymous namespace gives each
+    // translation unit its own unique definition.
+
     O << "\n#include <sycl/kernel_bundle.hpp>\n";
     O << "#include <sycl/detail/kernel_global_info.hpp>\n";
-    O << "namespace sycl {\n";
-    O << "inline namespace _V1 {\n";
-    O << "namespace detail {\n";
+    O << "namespace {\n";
     O << "struct GlobalMapUpdater {\n";
     O << "  GlobalMapUpdater() {\n";
     O << "    sycl::detail::free_function_info_map::add("
       << "sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, "
       << KernelDescs.size() << ");\n";
     O << "  }\n";
+    O << "  ~GlobalMapUpdater() {\n";
+    O << "    sycl::detail::free_function_info_map::remove("
+      << "sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, "
+      << KernelDescs.size() << ");\n";
+    O << "  }\n";
     O << "};\n";
     O << "static GlobalMapUpdater updater;\n";
-    O << "} // namespace detail\n";
-    O << "} // namespace _V1\n";
-    O << "} // namespace sycl\n";
+    O << "} // namespace\n";
   }
 }
 
diff --git a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
index e660bd3070874..56db32875ee51 100644
--- a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
+++ b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
@@ -1087,15 +1087,14 @@ namespace Testing::Tests {
 
 // CHECK: #include <sycl/kernel_bundle.hpp>
 // CHECK-NEXT: #include <sycl/detail/kernel_global_info.hpp>
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT: inline namespace _V1 {
-// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: namespace {
 // CHECK-NEXT: struct GlobalMapUpdater {
 // CHECK-NEXT:   GlobalMapUpdater() {
 // CHECK-NEXT:     sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 28);
 // CHECK-NEXT:   }
+// CHECK-NEXT:   ~GlobalMapUpdater() {
+// CHECK-NEXT:     sycl::detail::free_function_info_map::remove(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 28);
+// CHECK-NEXT:   }
 // CHECK-NEXT: };
 // CHECK-NEXT: static GlobalMapUpdater updater;
-// CHECK-NEXT: } // namespace detail
-// CHECK-NEXT: } // namespace _V1
-// CHECK-NEXT: } // namespace sycl
+// CHECK-NEXT: }
diff --git a/clang/test/CodeGenSYCL/free_function_int_header.cpp b/clang/test/CodeGenSYCL/free_function_int_header.cpp
index 4fe7a761e98c6..d589c6150e2a4 100644
--- a/clang/test/CodeGenSYCL/free_function_int_header.cpp
+++ b/clang/test/CodeGenSYCL/free_function_int_header.cpp
@@ -1567,15 +1567,14 @@ void ff_24(int arg) {
 
 // CHECK: #include <sycl/kernel_bundle.hpp>
 // CHECK-NEXT: #include <sycl/detail/kernel_global_info.hpp>
-// CHECK-NEXT: namespace sycl {
-// CHECK-NEXT: inline namespace _V1 {
-// CHECK-NEXT: namespace detail {
+// CHECK-NEXT: namespace {
 // CHECK-NEXT: struct GlobalMapUpdater {
 // CHECK-NEXT:  GlobalMapUpdater() {
 // CHECK-NEXT:     sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 33);
 // CHECK-NEXT:   }
+// CHECK-NEXT:  ~GlobalMapUpdater() {
+// CHECK-NEXT:     sycl::detail::free_function_info_map::remove(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 33);
+// CHECK-NEXT:   }
 // CHECK-NEXT: };
 // CHECK-NEXT: static GlobalMapUpdater updater;
-// CHECK-NEXT: } // namespace detail
-// CHECK-NEXT: } // namespace _V1
-// CHECK-NEXT: } // namespace sycl
+// CHECK-NEXT: }
diff --git a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
index fae7be0120ba9..f59598cda3db0 100644
--- a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
+++ b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
@@ -75,15 +75,14 @@ int main(){
 
 // CHECK-NORTC: #include <sycl/kernel_bundle.hpp>
 // CHECK-NORTC-NEXT: #include <sycl/detail/kernel_global_info.hpp>
-// CHECK-NORTC-NEXT: namespace sycl {
-// CHECK-NORTC-NEXT: inline namespace _V1 {
-// CHECK-NORTC-NEXT: namespace detail {
+// CHECK-NORTC-NEXT: namespace {
 // CHECK-NORTC-NEXT: struct GlobalMapUpdater {
 // CHECK-NORTC-NEXT:   GlobalMapUpdater() {
 // CHECK-NORTC-NEXT:     sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 3);
 // CHECK-NORTC-NEXT:   }
+// CHECK-NORTC-NEXT:   ~GlobalMapUpdater() {
+// CHECK-NORTC-NEXT:     sycl::detail::free_function_info_map::remove(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 3);
+// CHECK-NORTC-NEXT:   }
 // CHECK-NORTC-NEXT: };
 // CHECK-NORTC-NEXT: static GlobalMapUpdater updater;
-// CHECK-NORTC-NEXT: } // namespace detail
-// CHECK-NORTC-NEXT: } // namespace _V1
-// CHECK-NORTC-NEXT: } // namespace sycl
+// CHECK-NORTC-NEXT: }
diff --git a/clang/test/SemaSYCL/Inputs/sycl/detail/kernel_global_info.hpp b/clang/test/SemaSYCL/Inputs/sycl/detail/kernel_global_info.hpp
index 9903e088c81cc..3b0af6d9bf49b 100644
--- a/clang/test/SemaSYCL/Inputs/sycl/detail/kernel_global_info.hpp
+++ b/clang/test/SemaSYCL/Inputs/sycl/detail/kernel_global_info.hpp
@@ -10,14 +10,11 @@
 
 #include <sycl/detail/export.hpp>
 
-namespace sycl {
-inline namespace _V1 {
-namespace detail {
+namespace {
 namespace free_function_info_map {
 
 __SYCL_EXPORT void add(const void *DeviceGlobalPtr, const char *UniqueId);
+__SYCL_EXPORT void remove(const void *DeviceGlobalPtr, const char *UniqueId);
 
 } // namespace free_function_info_map
-} // namespace detail
-} // namespace _V1
-} // namespace sycl
+}
diff --git a/sycl/include/sycl/detail/kernel_global_info.hpp b/sycl/include/sycl/detail/kernel_global_info.hpp
index ac3cd76b2d92a..fb90b92f548ad 100644
--- a/sycl/include/sycl/detail/kernel_global_info.hpp
+++ b/sycl/include/sycl/detail/kernel_global_info.hpp
@@ -18,6 +18,8 @@ namespace free_function_info_map {
 __SYCL_EXPORT void add(const char *const *UniqueId,
                        const unsigned *DeviceGlobalPtr, unsigned Size);
 
+__SYCL_EXPORT void remove(const char *const *UniqueId,
+                          const unsigned *DeviceGlobalPtr, unsigned Size);
 } // namespace free_function_info_map
 } // namespace detail
 } // namespace _V1
diff --git a/sycl/source/detail/kernel_global_info.cpp b/sycl/source/detail/kernel_global_info.cpp
index fbc811ecf1227..704a13f6c7da0 100644
--- a/sycl/source/detail/kernel_global_info.cpp
+++ b/sycl/source/detail/kernel_global_info.cpp
@@ -23,6 +23,17 @@ __SYCL_EXPORT void add(const char *const *FreeFunctionNames,
       std::move(GlobalInfoToCopy));
 }
 
+__SYCL_EXPORT void remove(const char *const *FreeFunctionNames,
+                          const unsigned *FreeFunctionNumArgs, unsigned Size) {
+  std::unordered_map<std::string_view, unsigned> GlobalInfoToCopy;
+  for (size_t i = 0; i < Size; ++i) {
+    GlobalInfoToCopy[std::string_view{FreeFunctionNames[i]}] =
+        FreeFunctionNumArgs[i];
+  }
+  detail::ProgramManager::getInstance().unRegisterKernelGlobalInfo(
+      std::move(GlobalInfoToCopy));
+}
+
 } // namespace detail::free_function_info_map
 } // namespace _V1
 } // namespace sycl
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 2ca4420dc0549..50f2582ce63fe 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -2418,12 +2418,22 @@ void ProgramManager::registerKernelGlobalInfo(
   if (m_FreeFunctionKernelGlobalInfo.empty())
     m_FreeFunctionKernelGlobalInfo = std::move(GlobalInfoToCopy);
   else {
-    for (auto &GlobalInfo : GlobalInfoToCopy) {
+    for (auto &GlobalInfo : GlobalInfoToCopy)
       m_FreeFunctionKernelGlobalInfo.insert(GlobalInfo);
-    }
   }
 }
 
+// Remove entries from m_FreeFunctionKernelGlobalInfo that matches
+// the ones in GlobalInfoToCopy. This function is called when a shared
+// library consisting of SYCL kernels is unloaded.
+void ProgramManager::unRegisterKernelGlobalInfo(
+    std::unordered_map<std::string_view, unsigned> &&GlobalInfoToCopy) {
+  std::lock_guard<std::mutex> Guard(MNativeProgramsMutex);
+
+  for (const auto &GlobalInfo : GlobalInfoToCopy)
+    m_FreeFunctionKernelGlobalInfo.erase(GlobalInfo.first);
+}
+
 std::optional<unsigned>
 ProgramManager::getKernelGlobalInfoDesc(const char *UniqueId) {
   std::lock_guard<std::mutex> Guard(MNativeProgramsMutex);
diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp
index b9d0dc700f77c..73eca2cd86e0a 100644
--- a/sycl/source/detail/program_manager/program_manager.hpp
+++ b/sycl/source/detail/program_manager/program_manager.hpp
@@ -259,6 +259,12 @@ class ProgramManager {
   void registerKernelGlobalInfo(
       std::unordered_map<std::string_view, unsigned> &&GlobalInfoToCopy);
 
+  // The function removes kernel global descriptors from the
+  // kernel global map when a shared library consisting of SYCL kernels
+  // is unloaded.
+  void unRegisterKernelGlobalInfo(
+      std::unordered_map<std::string_view, unsigned> &&GlobalInfoToCopy);
+
   // The function returns a pointer to the kernel global desc identified by
   // the unique ID from the kernel global map.
   std::optional<unsigned> getKernelGlobalInfoDesc(const char *UniqueId);
diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump
index bfb11f8c79fe8..e9fdf8dd5215e 100644
--- a/sycl/test/abi/sycl_symbols_linux.dump
+++ b/sycl/test/abi/sycl_symbols_linux.dump
@@ -3344,6 +3344,7 @@ _ZN4sycl3_V16detail21LocalAccessorBaseHostC1ENS0_5rangeILi3EEEiiRKNS0_13property
 _ZN4sycl3_V16detail21LocalAccessorBaseHostC2ENS0_5rangeILi3EEEiiRKNS0_13property_listE
 _ZN4sycl3_V16detail22addHostAccessorAndWaitEPNS1_16AccessorImplHostE
 _ZN4sycl3_V16detail22free_function_info_map3addEPKPKcPKjj
+_ZN4sycl3_V16detail22free_function_info_map6removeEPKPKcPKjj
 _ZN4sycl3_V16detail22get_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EENS0_12bundle_stateE
 _ZN4sycl3_V16detail22get_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EENS0_12bundle_stateERKSt8functionIFbRKSt10shared_ptrINS1_17device_image_implEEEE
 _ZN4sycl3_V16detail22get_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EERKNS0_4spanIcLm18446744073709551615EEENS0_12bundle_stateE
diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump
index b2c805ac1b9d0..aed080bc22446 100644
--- a/sycl/test/abi/sycl_symbols_windows.dump
+++ b/sycl/test/abi/sycl_symbols_windows.dump
@@ -4393,6 +4393,7 @@
 ?release_external_semaphore@experimental@oneapi@ext@_V1@sycl@@YAXUexternal_semaphore@12345@AEBVqueue@45@@Z
 ?release_from_device_copy@experimental@oneapi@ext@_V1@sycl@@YAXPEBXAEBVcontext@45@@Z
 ?release_from_device_copy@experimental@oneapi@ext@_V1@sycl@@YAXPEBXAEBVqueue@45@@Z
+?remove@free_function_info_map@detail@_V1@sycl@@YAXPEBQEBDPEBII@Z
 ?removeDuplicateDevices@detail@_V1@sycl@@YA?BV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@AEBV45@@Z
 ?remquo_impl@detail@_V1@sycl@@YA?AVhalf@half_impl@123@V45123@0PEAH@Z
 ?remquo_impl@detail@_V1@sycl@@YAMMMPEAH@Z