diff --git a/sycl/source/detail/config.cpp b/sycl/source/detail/config.cpp
index 63e20207962f9..663747396ad82 100644
--- a/sycl/source/detail/config.cpp
+++ b/sycl/source/detail/config.cpp
@@ -180,6 +180,34 @@ const std::array<std::pair<std::string, backend>, 8> &getSyclBeMap() {
        {"*", backend::all}}};
   return SyclBeMap;
 }
+namespace {
+
+unsigned int parseLevel(const char *ValStr) {
+  unsigned int intVal = 0;
+
+  if (ValStr) {
+    try {
+      intVal = std::stoul(ValStr);
+    } catch (...) {
+      // If the value is not null and not a number, it is considered
+      // to enable disk cache tracing. This is the legacy behavior.
+      intVal = 1;
+    }
+  }
+
+  // Legacy behavior.
+  if (intVal > 7)
+    intVal = 1;
+
+  return intVal;
+}
+
+} // namespace
+
+void SYCLConfigTrace::reset() { Level = parseLevel(BaseT::getRawValue()); }
+
+unsigned int SYCLConfigTrace::Level =
+    parseLevel(SYCLConfigTrace::BaseT::getRawValue());
 
 } // namespace detail
 } // namespace _V1
diff --git a/sycl/source/detail/config.hpp b/sycl/source/detail/config.hpp
index 2eb3716c76a05..79cb8ab6f242c 100644
--- a/sycl/source/detail/config.hpp
+++ b/sycl/source/detail/config.hpp
@@ -709,52 +709,19 @@ template <> class SYCLConfig<SYCL_JIT_AMDGCN_PTX_TARGET_FEATURES> {
 // tracing of the corresponding caches. If the input value is not null and
 // not a valid number, the disk cache tracing will be enabled (depreciated
 // behavior). The default value is 0 and no tracing is enabled.
-template <> class SYCLConfig<SYCL_CACHE_TRACE> {
+class SYCLConfigTrace {
   using BaseT = SYCLConfigBase<SYCL_CACHE_TRACE>;
   enum TraceBitmask { DiskCache = 1, InMemCache = 2, KernelCompiler = 4 };
 
 public:
-  static unsigned int get() { return getCachedValue(); }
-  static void reset() { (void)getCachedValue(true); }
-  static bool isTraceDiskCache() {
-    return getCachedValue() & TraceBitmask::DiskCache;
-  }
-  static bool isTraceInMemCache() {
-    return getCachedValue() & TraceBitmask::InMemCache;
-  }
-  static bool isTraceKernelCompiler() {
-    return getCachedValue() & TraceBitmask::KernelCompiler;
-  }
+  static unsigned int get() { return Level; }
+  static void reset();
+  static bool isTraceDiskCache() { return Level & DiskCache; }
+  static bool isTraceInMemCache() { return Level & InMemCache; }
+  static bool isTraceKernelCompiler() { return Level & KernelCompiler; }
 
 private:
-  static unsigned int getCachedValue(bool ResetCache = false) {
-    const auto Parser = []() {
-      const char *ValStr = BaseT::getRawValue();
-      int intVal = 0;
-
-      if (ValStr) {
-        try {
-          intVal = std::stoi(ValStr);
-        } catch (...) {
-          // If the value is not null and not a number, it is considered
-          // to enable disk cache tracing. This is the legacy behavior.
-          intVal = 1;
-        }
-      }
-
-      // Legacy behavior.
-      if (intVal > 7)
-        intVal = 1;
-
-      return intVal;
-    };
-
-    static unsigned int Level = Parser();
-    if (ResetCache)
-      Level = Parser();
-
-    return Level;
-  }
+  static unsigned int Level;
 };
 
 // SYCL_IN_MEM_CACHE_EVICTION_THRESHOLD accepts an integer that specifies
diff --git a/sycl/source/detail/context_impl.cpp b/sycl/source/detail/context_impl.cpp
index 6fb2dd375fe37..026d1289e05ca 100644
--- a/sycl/source/detail/context_impl.cpp
+++ b/sycl/source/detail/context_impl.cpp
@@ -263,22 +263,6 @@ context_impl::get_backend_info<info::device::backend_version>() const {
 }
 #endif
 
-ur_context_handle_t &context_impl::getHandleRef() { return MContext; }
-const ur_context_handle_t &context_impl::getHandleRef() const {
-  return MContext;
-}
-
-KernelProgramCache &context_impl::getKernelProgramCache() const {
-  return MKernelProgramCache;
-}
-
-bool context_impl::hasDevice(const detail::device_impl &Device) const {
-  for (device_impl *D : MDevices)
-    if (D == &Device)
-      return true;
-  return false;
-}
-
 device_impl *
 context_impl::findMatchingDeviceImpl(ur_device_handle_t &DeviceUR) const {
   for (device_impl *D : MDevices)
diff --git a/sycl/source/detail/context_impl.hpp b/sycl/source/detail/context_impl.hpp
index 0ae3df8dcf397..5cf5c1f597830 100644
--- a/sycl/source/detail/context_impl.hpp
+++ b/sycl/source/detail/context_impl.hpp
@@ -116,7 +116,7 @@ class context_impl : public std::enable_shared_from_this<context_impl> {
   /// reference will be invalid if context_impl was destroyed.
   ///
   /// \return an instance of raw UR context handle.
-  ur_context_handle_t &getHandleRef();
+  ur_context_handle_t &getHandleRef() { return MContext; }
 
   /// Gets the underlying context object (if any) without reference count
   /// modification.
@@ -126,7 +126,7 @@ class context_impl : public std::enable_shared_from_this<context_impl> {
   /// reference will be invalid if context_impl was destroyed.
   ///
   /// \return an instance of raw UR context handle.
-  const ur_context_handle_t &getHandleRef() const;
+  const ur_context_handle_t &getHandleRef() const { return MContext; }
 
   devices_range getDevices() const { return MDevices; }
 
@@ -151,10 +151,17 @@ class context_impl : public std::enable_shared_from_this<context_impl> {
     return {MCachedLibPrograms, MCachedLibProgramsMutex};
   }
 
-  KernelProgramCache &getKernelProgramCache() const;
+  KernelProgramCache &getKernelProgramCache() const {
+    return MKernelProgramCache;
+  }
 
   /// Returns true if and only if context contains the given device.
-  bool hasDevice(const detail::device_impl &Device) const;
+  bool hasDevice(const detail::device_impl &Device) const {
+    for (device_impl *D : MDevices)
+      if (D == &Device)
+        return true;
+    return false;
+  }
 
   /// Returns true if and only if the device can be used within this context.
   /// For OpenCL this is currently equivalent to hasDevice, for other backends
diff --git a/sycl/source/detail/device_kernel_info.cpp b/sycl/source/detail/device_kernel_info.cpp
index 526f160c6596b..459e32fa4caf9 100644
--- a/sycl/source/detail/device_kernel_info.cpp
+++ b/sycl/source/detail/device_kernel_info.cpp
@@ -74,27 +74,6 @@ void DeviceKernelInfo::setCompileTimeInfoIfNeeded(
   assert(Info == *this);
 }
 
-FastKernelSubcacheT &DeviceKernelInfo::getKernelSubcache() {
-  assertInitialized();
-  return MFastKernelSubcache;
-}
-bool DeviceKernelInfo::usesAssert() {
-  assertInitialized();
-  return MUsesAssert;
-}
-const std::optional<int> &DeviceKernelInfo::getImplicitLocalArgPos() {
-  assertInitialized();
-  return MImplicitLocalArgPos;
-}
-
-bool DeviceKernelInfo::isCompileTimeInfoSet() const { return KernelSize != 0; }
-
-void DeviceKernelInfo::assertInitialized() {
-#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
-  assert(MInitialized.load() && "Data needs to be initialized before use");
-#endif
-}
-
 } // namespace detail
 } // namespace _V1
 } // namespace sycl
diff --git a/sycl/source/detail/device_kernel_info.hpp b/sycl/source/detail/device_kernel_info.hpp
index 0ea4ff2d051e6..5d64b18e97f75 100644
--- a/sycl/source/detail/device_kernel_info.hpp
+++ b/sycl/source/detail/device_kernel_info.hpp
@@ -107,13 +107,28 @@ class DeviceKernelInfo : public CompileTimeKernelInfoTy {
 #endif
   void setCompileTimeInfoIfNeeded(const CompileTimeKernelInfoTy &Info);
 
-  FastKernelSubcacheT &getKernelSubcache();
-  bool usesAssert();
-  const std::optional<int> &getImplicitLocalArgPos();
+  FastKernelSubcacheT &getKernelSubcache() {
+    assertInitialized();
+    return MFastKernelSubcache;
+  }
+
+  bool usesAssert() const {
+    assertInitialized();
+    return MUsesAssert;
+  }
+
+  std::optional<int> getImplicitLocalArgPos() const {
+    assertInitialized();
+    return MImplicitLocalArgPos;
+  }
 
 private:
-  void assertInitialized();
-  bool isCompileTimeInfoSet() const;
+  void assertInitialized() const {
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+    assert(MInitialized.load() && "Data needs to be initialized before use");
+#endif
+  }
+  bool isCompileTimeInfoSet() const { return KernelSize != 0; }
 
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   std::atomic<bool> MInitialized = false;
diff --git a/sycl/source/detail/global_handler.cpp b/sycl/source/detail/global_handler.cpp
index eb7d11d3b29d4..97e67523cff86 100644
--- a/sycl/source/detail/global_handler.cpp
+++ b/sycl/source/detail/global_handler.cpp
@@ -105,20 +105,7 @@ void GlobalHandler::TraceEventXPTI(const char *Message) {
 #endif
 }
 
-GlobalHandler *&GlobalHandler::getInstancePtr() {
-  static GlobalHandler *RTGlobalObjHandler = new GlobalHandler();
-  return RTGlobalObjHandler;
-}
-
-GlobalHandler &GlobalHandler::instance() {
-  GlobalHandler *RTGlobalObjHandler = GlobalHandler::getInstancePtr();
-  assert(RTGlobalObjHandler && "Handler must not be deallocated earlier");
-  return *RTGlobalObjHandler;
-}
-
-bool GlobalHandler::isInstanceAlive() {
-  return GlobalHandler::getInstancePtr();
-}
+GlobalHandler *GlobalHandler::RTGlobalObjHandler = new GlobalHandler();
 
 template <typename T, typename... Types>
 T &GlobalHandler::getOrCreate(InstWithLock<T> &IWL, Types &&...Args) {
@@ -331,8 +318,7 @@ void GlobalHandler::drainThreadPool() {
 //  2) when process is being terminated
 void shutdown_early(bool CanJoinThreads = true) {
   const LockGuard Lock{GlobalHandler::MSyclGlobalHandlerProtector};
-  GlobalHandler *&Handler = GlobalHandler::getInstancePtr();
-  if (!Handler)
+  if (!GlobalHandler::RTGlobalObjHandler)
     return;
 
 #if defined(XPTI_ENABLE_INSTRUMENTATION) && defined(_WIN32)
@@ -342,26 +328,26 @@ void shutdown_early(bool CanJoinThreads = true) {
 #endif
 
   // Now that we are shutting down, we will no longer defer MemObj releases.
-  Handler->endDeferredRelease();
+  GlobalHandler::RTGlobalObjHandler->endDeferredRelease();
 
   // Ensure neither host task is working so that no default context is accessed
   // upon its release
-  Handler->prepareSchedulerToRelease(true);
+  GlobalHandler::RTGlobalObjHandler->prepareSchedulerToRelease(true);
 
-  if (Handler->MHostTaskThreadPool.Inst) {
-    Handler->MHostTaskThreadPool.Inst->finishAndWait(CanJoinThreads);
-    Handler->MHostTaskThreadPool.Inst.reset(nullptr);
+  if (GlobalHandler::RTGlobalObjHandler->MHostTaskThreadPool.Inst) {
+    GlobalHandler::RTGlobalObjHandler->MHostTaskThreadPool.Inst->finishAndWait(
+        CanJoinThreads);
+    GlobalHandler::RTGlobalObjHandler->MHostTaskThreadPool.Inst.reset(nullptr);
   }
 
   // This releases OUR reference to the default context, but
   // other may yet have refs
-  Handler->releaseDefaultContexts();
+  GlobalHandler::RTGlobalObjHandler->releaseDefaultContexts();
 }
 
 void shutdown_late() {
   const LockGuard Lock{GlobalHandler::MSyclGlobalHandlerProtector};
-  GlobalHandler *&Handler = GlobalHandler::getInstancePtr();
-  if (!Handler)
+  if (!GlobalHandler::RTGlobalObjHandler)
     return;
 
 #if defined(XPTI_ENABLE_INSTRUMENTATION) && defined(_WIN32)
@@ -371,26 +357,27 @@ void shutdown_late() {
 #endif
 
   // First, release resources, that may access adapters.
-  Handler->MPlatformCache.Inst.reset(nullptr);
-  Handler->MScheduler.Inst.reset(nullptr);
-  Handler->MProgramManager.Inst.reset(nullptr);
+  GlobalHandler::RTGlobalObjHandler->MPlatformCache.Inst.reset(nullptr);
+  GlobalHandler::RTGlobalObjHandler->MScheduler.Inst.reset(nullptr);
+  GlobalHandler::RTGlobalObjHandler->MProgramManager.Inst.reset(nullptr);
 
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   // Kernel cache, which is part of device kernel info,
   // stores handles to the adapter, so clear it before releasing adapters.
-  Handler->MDeviceKernelInfoStorage.Inst.reset(nullptr);
+  GlobalHandler::RTGlobalObjHandler->MDeviceKernelInfoStorage.Inst.reset(
+      nullptr);
 #endif
 
   // Clear the adapters and reset the instance if it was there.
-  Handler->unloadAdapters();
-  if (Handler->MAdapters.Inst)
-    Handler->MAdapters.Inst.reset(nullptr);
+  GlobalHandler::RTGlobalObjHandler->unloadAdapters();
+  if (GlobalHandler::RTGlobalObjHandler->MAdapters.Inst)
+    GlobalHandler::RTGlobalObjHandler->MAdapters.Inst.reset(nullptr);
 
-  Handler->MXPTIRegistry.Inst.reset(nullptr);
+  GlobalHandler::RTGlobalObjHandler->MXPTIRegistry.Inst.reset(nullptr);
 
   // Release the rest of global resources.
-  delete Handler;
-  Handler = nullptr;
+  delete GlobalHandler::RTGlobalObjHandler;
+  GlobalHandler::RTGlobalObjHandler = nullptr;
 }
 
 #ifdef _WIN32
diff --git a/sycl/source/detail/global_handler.hpp b/sycl/source/detail/global_handler.hpp
index ec7bf7da48b6a..5b0a01d8cef7d 100644
--- a/sycl/source/detail/global_handler.hpp
+++ b/sycl/source/detail/global_handler.hpp
@@ -48,14 +48,11 @@ class DeviceKernelInfo;
 /// construction or destruction is generated anyway.
 class GlobalHandler {
 public:
-  /// \return a reference to a GlobalHandler singleton instance. Memory for
-  /// storing objects is allocated on first call. The reference is valid as long
-  /// as runtime library is loaded (i.e. untill `DllMain` or
+  static bool isInstanceAlive() { return RTGlobalObjHandler != nullptr; }
+  /// \return a reference to a GlobalHandler singleton instance. The reference
+  /// is valid as long as runtime library is loaded (i.e. untill `DllMain` or
   /// `__attribute__((destructor))` is called).
-  static GlobalHandler &instance();
-
-  /// \return true if the instance has not been deallocated yet.
-  static bool isInstanceAlive();
+  static GlobalHandler &instance() { return *RTGlobalObjHandler; }
 
   GlobalHandler(const GlobalHandler &) = delete;
   GlobalHandler(GlobalHandler &&) = delete;
@@ -96,19 +93,18 @@ class GlobalHandler {
   void attachScheduler(Scheduler *Scheduler);
 
 private:
+  // Constructor and destructor are declared out-of-line to allow incomplete
+  // types as template arguments to unique_ptr.
+  GlobalHandler();
+  ~GlobalHandler();
+
   bool OkToDefer = true;
 
   friend void shutdown_early(bool);
   friend void shutdown_late();
   friend class ObjectUsageCounter;
-  static GlobalHandler *&getInstancePtr();
   static SpinLock MSyclGlobalHandlerProtector;
 
-  // Constructor and destructor are declared out-of-line to allow incomplete
-  // types as template arguments to unique_ptr.
-  GlobalHandler();
-  ~GlobalHandler();
-
   template <typename T> struct InstWithLock {
     std::unique_ptr<T> Inst;
     SpinLock Lock;
@@ -135,7 +131,10 @@ class GlobalHandler {
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   InstWithLock<std::deque<DeviceKernelInfo>> MDeviceKernelInfoStorage;
 #endif
+
+  static GlobalHandler *RTGlobalObjHandler;
 };
+
 } // namespace detail
 } // namespace _V1
 } // namespace sycl
diff --git a/sycl/source/detail/kernel_program_cache.cpp b/sycl/source/detail/kernel_program_cache.cpp
index e73c0a19ef589..e778e13e6ce91 100644
--- a/sycl/source/detail/kernel_program_cache.cpp
+++ b/sycl/source/detail/kernel_program_cache.cpp
@@ -12,6 +12,18 @@
 namespace sycl {
 inline namespace _V1 {
 namespace detail {
+
+void KernelProgramCache::traceKernelImpl(const char *Msg,
+                                         KernelNameStrRefT KernelName,
+                                         bool IsFastKernelCache) {
+  std::string Identifier =
+      "[IsFastCache: " + std::to_string(IsFastKernelCache) +
+      "][Key:{Name = " + KernelName.data() + "}]: ";
+
+  std::cerr << "[In-Memory Cache][Thread Id:" << std::this_thread::get_id()
+            << "][Kernel Cache]" << Identifier << Msg << std::endl;
+}
+
 adapter_impl &KernelProgramCache::getAdapter() {
   return MParentContext.getAdapter();
 }
diff --git a/sycl/source/detail/kernel_program_cache.hpp b/sycl/source/detail/kernel_program_cache.hpp
index d0a6b398528c4..036a91e141c41 100644
--- a/sycl/source/detail/kernel_program_cache.hpp
+++ b/sycl/source/detail/kernel_program_cache.hpp
@@ -331,7 +331,7 @@ class KernelProgramCache {
   template <typename MsgType>
   static inline void traceProgram(const MsgType &Msg,
                                   const ProgramCacheKeyT &CacheKey) {
-    if (!SYCLConfig<SYCL_CACHE_TRACE>::isTraceInMemCache())
+    if (!SYCLConfigTrace::isTraceInMemCache())
       return;
 
     int ImageId = CacheKey.first.second;
@@ -361,21 +361,15 @@ class KernelProgramCache {
               << "][Program Cache]" << Identifier << Msg << std::endl;
   }
 
+  static void traceKernelImpl(const char *Msg, KernelNameStrRefT KernelName,
+                              bool IsFastKernelCache);
+
   // Sends message to std:cerr stream when SYCL_CACHE_TRACE environemnt is
   // set.
-  template <typename MsgType>
-  static inline void traceKernel(const MsgType &Msg,
-                                 KernelNameStrRefT KernelName,
-                                 bool IsFastKernelCache = false) {
-    if (!SYCLConfig<SYCL_CACHE_TRACE>::isTraceInMemCache())
-      return;
-
-    std::string Identifier =
-        "[IsFastCache: " + std::to_string(IsFastKernelCache) +
-        "][Key:{Name = " + KernelName.data() + "}]: ";
-
-    std::cerr << "[In-Memory Cache][Thread Id:" << std::this_thread::get_id()
-              << "][Kernel Cache]" << Identifier << Msg << std::endl;
+  static void traceKernel(const char *Msg, KernelNameStrRefT KernelName,
+                          bool isFastKernelCache = false) {
+    if (__builtin_expect(SYCLConfigTrace::isTraceInMemCache(), false))
+      traceKernelImpl(Msg, KernelName, isFastKernelCache);
   }
 
   Locked<ProgramCache> acquireCachedPrograms() {
@@ -513,7 +507,7 @@ class KernelProgramCache {
         auto LockedCacheKP = acquireKernelsPerProgramCache();
         // List kernels that are to be removed from the cache, if tracing is
         // enabled.
-        if (SYCLConfig<SYCL_CACHE_TRACE>::isTraceInMemCache()) {
+        if (SYCLConfigTrace::isTraceInMemCache()) {
           for (const auto &Kernel : LockedCacheKP.get()[NativePrg])
             traceKernel("Kernel evicted.", Kernel.first);
         }
diff --git a/sycl/source/detail/persistent_device_code_cache.hpp b/sycl/source/detail/persistent_device_code_cache.hpp
index 5d73db711d0cc..7747a415a282b 100644
--- a/sycl/source/detail/persistent_device_code_cache.hpp
+++ b/sycl/source/detail/persistent_device_code_cache.hpp
@@ -222,8 +222,7 @@ class PersistentDeviceCodeCache {
 
   /* Sends message to std:cerr stream when SYCL_CACHE_TRACE environemnt is set*/
   static void trace(const std::string &msg, const std::string &path = "") {
-    static const bool traceEnabled =
-        SYCLConfig<SYCL_CACHE_TRACE>::isTraceDiskCache();
+    static const bool traceEnabled = SYCLConfigTrace::isTraceDiskCache();
     if (traceEnabled) {
       auto outputPath = path;
       std::replace(outputPath.begin(), outputPath.end(), '\\', '/');
@@ -232,8 +231,7 @@ class PersistentDeviceCodeCache {
   }
   static void trace_KernelCompiler(const std::string &msg,
                                    const std::string &path = "") {
-    static const bool traceEnabled =
-        SYCLConfig<SYCL_CACHE_TRACE>::isTraceKernelCompiler();
+    static const bool traceEnabled = SYCLConfigTrace::isTraceKernelCompiler();
     if (traceEnabled) {
       auto outputPath = path;
       std::replace(outputPath.begin(), outputPath.end(), '\\', '/');
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 1d39a89a4dd45..2eac9c08d7f73 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -66,10 +66,6 @@ static void enableITTAnnotationsIfNeeded(const ur_program_handle_t &Prog,
   }
 }
 
-ProgramManager &ProgramManager::getInstance() {
-  return GlobalHandler::instance().getProgramManager();
-}
-
 static Managed<ur_program_handle_t>
 createBinaryProgram(context_impl &Context, devices_range Devices,
                     const uint8_t **Binaries, size_t *Lengths,
@@ -1813,14 +1809,6 @@ void ProgramManager::cacheKernelImplicitLocalArg(
     }
 }
 
-std::optional<int>
-ProgramManager::kernelImplicitLocalArgPos(KernelNameStrRefT KernelName) const {
-  auto it = m_KernelImplicitLocalArgPos.find(KernelName);
-  if (it != m_KernelImplicitLocalArgPos.end())
-    return it->second;
-  return {};
-}
-
 DeviceKernelInfo &ProgramManager::getOrCreateDeviceKernelInfo(
     const CompileTimeKernelInfoTy &Info) {
   std::lock_guard<std::mutex> Guard(m_DeviceKernelInfoMapMutex);
@@ -2355,24 +2343,6 @@ ProgramManager::getBinImageState(const RTDeviceBinaryImage *BinImage) {
                                                 : sycl::bundle_state::object;
 }
 
-std::optional<kernel_id>
-ProgramManager::tryGetSYCLKernelID(KernelNameStrRefT KernelName) {
-  std::lock_guard<std::mutex> KernelIDsGuard(m_KernelIDsMutex);
-
-  auto KernelID = m_KernelName2KernelIDs.find(KernelName);
-  if (KernelID == m_KernelName2KernelIDs.end())
-    return std::nullopt;
-
-  return KernelID->second;
-}
-
-kernel_id ProgramManager::getSYCLKernelID(KernelNameStrRefT KernelName) {
-  if (std::optional<kernel_id> MaybeKernelID = tryGetSYCLKernelID(KernelName))
-    return *MaybeKernelID;
-  throw exception(make_error_code(errc::runtime),
-                  "No kernel found with the specified name");
-}
-
 bool ProgramManager::hasCompatibleImage(const device_impl &DeviceImpl) {
   std::lock_guard<std::mutex> Guard(m_KernelIDsMutex);
 
diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp
index 73eca2cd86e0a..c3864edebde14 100644
--- a/sycl/source/detail/program_manager/program_manager.hpp
+++ b/sycl/source/detail/program_manager/program_manager.hpp
@@ -133,7 +133,9 @@ class ProgramManager {
 public:
   // Returns the single instance of the program manager for the entire
   // process. Can only be called after staticInit is done.
-  static ProgramManager &getInstance();
+  static ProgramManager &getInstance() {
+    return GlobalHandler::instance().getProgramManager();
+  }
 
   const RTDeviceBinaryImage &getDeviceImage(KernelNameStrRefT KernelName,
                                             context_impl &ContextImpl,
@@ -571,6 +573,33 @@ class ProgramManager {
 
   friend class ::ProgramManagerTest;
 };
+
+inline std::optional<int>
+ProgramManager::kernelImplicitLocalArgPos(KernelNameStrRefT KernelName) const {
+  auto it = m_KernelImplicitLocalArgPos.find(KernelName);
+  if (it != m_KernelImplicitLocalArgPos.end())
+    return it->second;
+  return {};
+}
+
+inline std::optional<kernel_id>
+ProgramManager::tryGetSYCLKernelID(KernelNameStrRefT KernelName) {
+  std::lock_guard<std::mutex> KernelIDsGuard(m_KernelIDsMutex);
+
+  auto KernelID = m_KernelName2KernelIDs.find(KernelName);
+  if (KernelID == m_KernelName2KernelIDs.end())
+    return std::nullopt;
+
+  return KernelID->second;
+}
+
+inline kernel_id ProgramManager::getSYCLKernelID(KernelNameStrRefT KernelName) {
+  if (std::optional<kernel_id> MaybeKernelID = tryGetSYCLKernelID(KernelName))
+    return *MaybeKernelID;
+  throw exception(make_error_code(errc::runtime),
+                  "No kernel found with the specified name");
+}
+
 } // namespace detail
 } // namespace _V1
 } // namespace sycl
diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp
index 2285bbce42761..6586a92ae20b9 100644
--- a/sycl/source/detail/scheduler/commands.cpp
+++ b/sycl/source/detail/scheduler/commands.cpp
@@ -2279,23 +2279,6 @@ static void adjustNDRangePerKernel(NDRDescT &NDR, ur_kernel_handle_t Kernel,
   }
 }
 
-// We have the following mapping between dimensions with SPIR-V builtins:
-// 1D: id[0] -> x
-// 2D: id[0] -> y, id[1] -> x
-// 3D: id[0] -> z, id[1] -> y, id[2] -> x
-// So in order to ensure the correctness we update all the kernel
-// parameters accordingly.
-// Initially we keep the order of NDRDescT as it provided by the user, this
-// simplifies overall handling and do the reverse only when
-// the kernel is enqueued.
-void ReverseRangeDimensionsForKernel(NDRDescT &NDR) {
-  if (NDR.Dims > 1) {
-    std::swap(NDR.GlobalSize[0], NDR.GlobalSize[NDR.Dims - 1]);
-    std::swap(NDR.LocalSize[0], NDR.LocalSize[NDR.Dims - 1]);
-    std::swap(NDR.GlobalOffset[0], NDR.GlobalOffset[NDR.Dims - 1]);
-  }
-}
-
 ur_mem_flags_t AccessModeToUr(access::mode AccessorMode) {
   switch (AccessorMode) {
   case access::mode::read:
diff --git a/sycl/source/detail/scheduler/commands.hpp b/sycl/source/detail/scheduler/commands.hpp
index d47a5d9d9131f..1f3264c36a997 100644
--- a/sycl/source/detail/scheduler/commands.hpp
+++ b/sycl/source/detail/scheduler/commands.hpp
@@ -796,7 +796,22 @@ void applyFuncOnFilteredArgs(
   }
 }
 
-void ReverseRangeDimensionsForKernel(NDRDescT &NDR);
+// We have the following mapping between dimensions with SPIR-V builtins:
+// 1D: id[0] -> x
+// 2D: id[0] -> y, id[1] -> x
+// 3D: id[0] -> z, id[1] -> y, id[2] -> x
+// So in order to ensure the correctness we update all the kernel
+// parameters accordingly.
+// Initially we keep the order of NDRDescT as it provided by the user, this
+// simplifies overall handling and do the reverse only when
+// the kernel is enqueued.
+inline void ReverseRangeDimensionsForKernel(NDRDescT &NDR) {
+  if (NDR.Dims > 1) {
+    std::swap(NDR.GlobalSize[0], NDR.GlobalSize[NDR.Dims - 1]);
+    std::swap(NDR.LocalSize[0], NDR.LocalSize[NDR.Dims - 1]);
+    std::swap(NDR.GlobalOffset[0], NDR.GlobalOffset[NDR.Dims - 1]);
+  }
+}
 
 } // namespace detail
 } // namespace _V1
diff --git a/sycl/unittests/config/ConfigTests.cpp b/sycl/unittests/config/ConfigTests.cpp
index 2391bb608a61e..b3d0baeeced80 100644
--- a/sycl/unittests/config/ConfigTests.cpp
+++ b/sycl/unittests/config/ConfigTests.cpp
@@ -28,8 +28,8 @@ TEST(ConfigTests, CheckConfigProcessing) {
     File.close();
   }
   try {
-    sycl::detail::SYCLConfig<sycl::detail::SYCL_DEVICE_ALLOWLIST>::get();
-    throw std::logic_error("sycl::exception didn't throw");
+    sycl::detail::readConfig(true);
+    throw std::logic_error("sycl::exception didn't throw 1");
   } catch (sycl::exception &e) {
     EXPECT_EQ(
         std::string(
@@ -46,8 +46,8 @@ TEST(ConfigTests, CheckConfigProcessing) {
     File.close();
   }
   try {
-    sycl::detail::SYCLConfig<sycl::detail::SYCL_DEVICE_ALLOWLIST>::get();
-    throw std::logic_error("sycl::exception didn't throw");
+    sycl::detail::readConfig(true);
+    throw std::logic_error("sycl::exception didn't throw 2");
   } catch (sycl::exception &e) {
     EXPECT_EQ(
         std::string(
@@ -64,8 +64,8 @@ TEST(ConfigTests, CheckConfigProcessing) {
     File.close();
   }
   try {
-    sycl::detail::SYCLConfig<sycl::detail::SYCL_DEVICE_ALLOWLIST>::get();
-    throw std::logic_error("sycl::exception didn't throw");
+    sycl::detail::readConfig(true);
+    throw std::logic_error("sycl::exception didn't throw 3");
   } catch (sycl::exception &e) {
     EXPECT_EQ(
         std::string(
@@ -82,8 +82,8 @@ TEST(ConfigTests, CheckConfigProcessing) {
     File.close();
   }
   try {
-    sycl::detail::SYCLConfig<sycl::detail::SYCL_DEVICE_ALLOWLIST>::get();
-    throw std::logic_error("sycl::exception didn't throw");
+    sycl::detail::readConfig(true);
+    throw std::logic_error("sycl::exception didn't throw 4");
   } catch (sycl::exception &e) {
     EXPECT_EQ(
         std::string(
@@ -103,8 +103,8 @@ TEST(ConfigTests, CheckConfigProcessing) {
     File.close();
   }
   try {
-    sycl::detail::SYCLConfig<sycl::detail::SYCL_DEVICE_ALLOWLIST>::get();
-    throw std::logic_error("sycl::exception didn't throw");
+    sycl::detail::readConfig(true);
+    throw std::logic_error("sycl::exception didn't throw 5");
   } catch (sycl::exception &e) {
     EXPECT_TRUE(std::regex_match(
         e.what(),
@@ -121,8 +121,8 @@ TEST(ConfigTests, CheckConfigProcessing) {
     File.close();
   }
   try {
-    sycl::detail::SYCLConfig<sycl::detail::SYCL_DEVICE_ALLOWLIST>::get();
-    throw std::logic_error("sycl::exception didn't throw");
+    sycl::detail::readConfig(true);
+    throw std::logic_error("sycl::exception didn't throw 6");
   } catch (sycl::exception &e) {
     EXPECT_TRUE(std::regex_match(
         e.what(), std::regex("Variable name is more than ([\\d]+) or less "
@@ -142,8 +142,8 @@ TEST(ConfigTests, CheckConfigProcessing) {
     File.close();
   }
   try {
-    sycl::detail::SYCLConfig<sycl::detail::SYCL_DEVICE_ALLOWLIST>::get();
-    throw std::logic_error("sycl::exception didn't throw");
+    sycl::detail::readConfig(true);
+    throw std::logic_error("sycl::exception didn't throw 7");
   } catch (sycl::exception &e) {
     EXPECT_TRUE(std::regex_match(
         e.what(), std::regex("The value contains more than ([\\d]+) characters "
@@ -159,8 +159,8 @@ TEST(ConfigTests, CheckConfigProcessing) {
     File.close();
   }
   try {
-    sycl::detail::SYCLConfig<sycl::detail::SYCL_DEVICE_ALLOWLIST>::get();
-    throw std::logic_error("sycl::exception didn't throw");
+    sycl::detail::readConfig(true);
+    throw std::logic_error("sycl::exception didn't throw 8");
   } catch (sycl::exception &e) {
     EXPECT_TRUE(std::regex_match(
         e.what(), std::regex("The value contains more than ([\\d]+) characters "
@@ -176,8 +176,8 @@ TEST(ConfigTests, CheckConfigProcessing) {
     File.close();
   }
   try {
-    sycl::detail::SYCLConfig<sycl::detail::SYCL_DEVICE_ALLOWLIST>::get();
-    throw std::logic_error("sycl::exception didn't throw");
+    sycl::detail::readConfig(true);
+    throw std::logic_error("sycl::exception didn't throw 9");
   } catch (sycl::exception &e) {
     EXPECT_TRUE(std::regex_match(
         e.what(), std::regex("The value contains more than ([\\d]+) characters "
@@ -249,20 +249,17 @@ TEST(ConfigTests, CheckSyclCacheTraceTest) {
   // Lambda to test parsing of SYCL_CACHE_TRACE
   auto TestConfig = [](int expectedValue, int expectedDiskCache,
                        int expectedInMemCache, int expectedKernelCompiler) {
-    EXPECT_EQ(static_cast<unsigned int>(expectedValue),
-              SYCLConfig<SYCL_CACHE_TRACE>::get());
+    EXPECT_EQ(static_cast<unsigned int>(expectedValue), SYCLConfigTrace::get());
 
     EXPECT_EQ(
         expectedDiskCache,
-        static_cast<int>(
-            sycl::detail::SYCLConfig<SYCL_CACHE_TRACE>::isTraceDiskCache()));
+        static_cast<int>(sycl::detail::SYCLConfigTrace::isTraceDiskCache()));
     EXPECT_EQ(
         expectedInMemCache,
-        static_cast<int>(
-            sycl::detail::SYCLConfig<SYCL_CACHE_TRACE>::isTraceInMemCache()));
+        static_cast<int>(sycl::detail::SYCLConfigTrace::isTraceInMemCache()));
     EXPECT_EQ(expectedKernelCompiler,
-              static_cast<int>(sycl::detail::SYCLConfig<
-                               SYCL_CACHE_TRACE>::isTraceKernelCompiler()));
+              static_cast<int>(
+                  sycl::detail::SYCLConfigTrace::isTraceKernelCompiler()));
   };
 
   // Lambda to set SYCL_CACHE_TRACE
@@ -279,40 +276,40 @@ TEST(ConfigTests, CheckSyclCacheTraceTest) {
   TestConfig(0, 0, 0, 0);
 
   SetSyclCacheTraceEnv("1");
-  sycl::detail::SYCLConfig<SYCL_CACHE_TRACE>::reset();
+  sycl::detail::SYCLConfigTrace::reset();
   TestConfig(1, 1, 0, 0);
 
   SetSyclCacheTraceEnv("2");
-  sycl::detail::SYCLConfig<SYCL_CACHE_TRACE>::reset();
+  sycl::detail::SYCLConfigTrace::reset();
   TestConfig(2, 0, 1, 0);
 
   SetSyclCacheTraceEnv("3");
-  sycl::detail::SYCLConfig<SYCL_CACHE_TRACE>::reset();
+  sycl::detail::SYCLConfigTrace::reset();
   TestConfig(3, 1, 1, 0);
 
   SetSyclCacheTraceEnv("4");
-  sycl::detail::SYCLConfig<SYCL_CACHE_TRACE>::reset();
+  sycl::detail::SYCLConfigTrace::reset();
   TestConfig(4, 0, 0, 1);
 
   SetSyclCacheTraceEnv("5");
-  sycl::detail::SYCLConfig<SYCL_CACHE_TRACE>::reset();
+  sycl::detail::SYCLConfigTrace::reset();
   TestConfig(5, 1, 0, 1);
 
   SetSyclCacheTraceEnv("6");
-  sycl::detail::SYCLConfig<SYCL_CACHE_TRACE>::reset();
+  sycl::detail::SYCLConfigTrace::reset();
   TestConfig(6, 0, 1, 1);
 
   SetSyclCacheTraceEnv("7");
-  sycl::detail::SYCLConfig<SYCL_CACHE_TRACE>::reset();
+  sycl::detail::SYCLConfigTrace::reset();
   TestConfig(7, 1, 1, 1);
 
   SetSyclCacheTraceEnv("8");
-  sycl::detail::SYCLConfig<SYCL_CACHE_TRACE>::reset();
+  sycl::detail::SYCLConfigTrace::reset();
   TestConfig(1, 1, 0, 0);
 
   // Set random non-null value. It should default to 1.
   SetSyclCacheTraceEnv("random");
-  sycl::detail::SYCLConfig<SYCL_CACHE_TRACE>::reset();
+  sycl::detail::SYCLConfigTrace::reset();
   TestConfig(1, 1, 0, 0);
 
   // When SYCL_CACHE_TRACE is not set, it should default to 0.
@@ -321,7 +318,7 @@ TEST(ConfigTests, CheckSyclCacheTraceTest) {
 #else
   unsetenv("SYCL_CACHE_TRACE");
 #endif
-  sycl::detail::SYCLConfig<SYCL_CACHE_TRACE>::reset();
+  sycl::detail::SYCLConfigTrace::reset();
   TestConfig(0, 0, 0, 0);
 }