diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp
index e688af8f3190c..fdca59637d5c0 100644
--- a/sycl/source/detail/device_image_impl.hpp
+++ b/sycl/source/detail/device_image_impl.hpp
@@ -1354,6 +1354,7 @@ class device_image_impl
 using device_images_iterator =
     variadic_iterator<device_image_plain,
                       std::vector<device_image_plain>::const_iterator,
+                      const device_image_plain *,
                       std::set<device_image_impl *>::const_iterator>;
 class device_images_range : public iterator_range<device_images_iterator> {
 private:
@@ -1361,6 +1362,21 @@ class device_images_range : public iterator_range<device_images_iterator> {
 
 public:
   using Base::Base;
+  template <typename Container>
+  decltype(std::declval<Base>().to<Container>()) to() const {
+    return this->Base::to<Container>();
+  }
+
+  template <typename Container>
+  std::enable_if_t<std::is_same_v<Container, std::vector<ur_program_handle_t>>,
+                   Container>
+  to() const {
+    std::vector<ur_program_handle_t> ProgramHandles;
+    ProgramHandles.reserve(size());
+    std::transform(begin(), end(), std::back_inserter(ProgramHandles),
+                   [](device_image_impl &Img) { return Img.get_ur_program(); });
+    return ProgramHandles;
+  }
 };
 
 } // namespace detail
diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp
index bba35438c7495..7ee2490eaa26c 100644
--- a/sycl/source/detail/kernel_bundle_impl.hpp
+++ b/sycl/source/detail/kernel_bundle_impl.hpp
@@ -114,12 +114,11 @@ CreateLinkGraph(const std::vector<device_image_plain> &DevImages) {
 }
 
 inline void
-ThrowIfConflictingKernels(const std::vector<device_image_plain> &DevImages) {
+ThrowIfConflictingKernels(device_images_range DevImages) {
   std::set<std::string_view, std::less<>> SeenKernelNames;
   std::set<std::string_view, std::less<>> Conflicts;
-  for (const device_image_plain &DevImage : DevImages) {
-    const KernelNameSetT &KernelNames =
-        getSyclObjImpl(DevImage)->getKernelNames();
+  for (const device_image_impl &DevImage : DevImages) {
+    const KernelNameSetT &KernelNames = DevImage.getKernelNames();
     std::vector<std::string_view> Intersect;
     std::set_intersection(SeenKernelNames.begin(), SeenKernelNames.end(),
                           KernelNames.begin(), KernelNames.end(),
@@ -513,9 +512,7 @@ class kernel_bundle_impl
         // In dynamic linking, AOT binaries count as results as well.
         LinkedResults.insert(LinkedResults.end(), AOTImgs.begin(),
                              AOTImgs.end());
-        sycl::span<device_image_plain, dynamic_extent> LinkedResultsSpan(
-            LinkedResults.data(), LinkedResults.size());
-        detail::ProgramManager::getInstance().dynamicLink(LinkedResultsSpan);
+        detail::ProgramManager::getInstance().dynamicLink(LinkedResults);
       }
 
       MDeviceImages.insert(MDeviceImages.end(), LinkedResults.begin(),
@@ -536,12 +533,9 @@ class kernel_bundle_impl
                   }))
         continue;
 
-      const std::vector<device_image_plain> &AllDevImgs =
-          DeviceImageWithDeps->getAll();
-      sycl::span<const device_image_plain> AllDevImgsSpan(AllDevImgs);
       std::vector<device_image_plain> LinkedResults =
-          detail::ProgramManager::getInstance().link(AllDevImgsSpan, MDevices,
-                                                     PropList);
+          detail::ProgramManager::getInstance().link(
+              DeviceImageWithDeps->getAll(), MDevices, PropList);
       MDeviceImages.insert(MDeviceImages.end(), LinkedResults.begin(),
                            LinkedResults.end());
       MUniqueDeviceImages.insert(MUniqueDeviceImages.end(),
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 278b2395839ad..ead96795b912c 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -2919,24 +2919,22 @@ ProgramManager::compile(const DevImgPlainWithDeps &ImgWithDeps,
 // Returns a merged device binary image, new set of kernel IDs and new
 // specialization constant data.
 static const RTDeviceBinaryImage *
-mergeImageData(sycl::span<const device_image_plain> Imgs,
-               std::vector<kernel_id> &KernelIDs,
+mergeImageData(device_images_range Imgs, std::vector<kernel_id> &KernelIDs,
                std::vector<unsigned char> &NewSpecConstBlob,
                device_image_impl::SpecConstMapT &NewSpecConstMap,
                std::unique_ptr<DynRTDeviceBinaryImage> &MergedImageStorage) {
-  for (const device_image_plain &Img : Imgs) {
-    device_image_impl &DeviceImageImpl = *getSyclObjImpl(Img);
+  for (device_image_impl &Img : Imgs) {
     // Duplicates are not expected here, otherwise urProgramLink should fail
-    KernelIDs.insert(KernelIDs.end(), DeviceImageImpl.get_kernel_ids().begin(),
-                     DeviceImageImpl.get_kernel_ids().end());
+    KernelIDs.insert(KernelIDs.end(), Img.get_kernel_ids().begin(),
+                     Img.get_kernel_ids().end());
     // To be able to answer queries about specialziation constants, the new
     // device image should have the specialization constants from all the linked
     // images.
     const std::lock_guard<std::mutex> SpecConstLock(
-        DeviceImageImpl.get_spec_const_data_lock());
+        Img.get_spec_const_data_lock());
     // Copy all map entries to the new map. Since the blob will be copied to
     // the end of the new blob we need to move the blob offset of each entry.
-    for (const auto &SpecConstIt : DeviceImageImpl.get_spec_const_data_ref()) {
+    for (const auto &SpecConstIt : Img.get_spec_const_data_ref()) {
       std::vector<device_image_impl::SpecConstDescT> &NewDescEntries =
           NewSpecConstMap[SpecConstIt.first];
 
@@ -2954,21 +2952,21 @@ mergeImageData(sycl::span<const device_image_plain> Imgs,
     // Copy the blob from the device image into the new blob. This moves the
     // offsets of the following blobs.
     NewSpecConstBlob.insert(NewSpecConstBlob.end(),
-                            DeviceImageImpl.get_spec_const_blob_ref().begin(),
-                            DeviceImageImpl.get_spec_const_blob_ref().end());
+                            Img.get_spec_const_blob_ref().begin(),
+                            Img.get_spec_const_blob_ref().end());
   }
   // device_image_impl expects kernel ids to be sorted for fast search
   std::sort(KernelIDs.begin(), KernelIDs.end(), LessByHash<kernel_id>{});
 
   // If there is only a single image, use it as the result.
   if (Imgs.size() == 1)
-    return getSyclObjImpl(Imgs[0])->get_bin_image_ref();
+    return Imgs.front().get_bin_image_ref();
 
   // Otherwise we create a dynamic image with the merged information.
   std::vector<const RTDeviceBinaryImage *> BinImgs;
   BinImgs.reserve(Imgs.size());
-  for (const device_image_plain &Img : Imgs) {
-    auto ImgBinRef = getSyclObjImpl(Img)->get_bin_image_ref();
+  for (device_image_impl &Img : Imgs) {
+    auto ImgBinRef = Img.get_bin_image_ref();
     // For some cases, like SYCL kernel compiler binaries, we don't have
     // binaries. For these we assume no properties associated, so they can be
     // safely ignored.
@@ -2980,8 +2978,8 @@ mergeImageData(sycl::span<const device_image_plain> Imgs,
 }
 
 std::vector<device_image_plain>
-ProgramManager::link(sycl::span<const device_image_plain> Imgs,
-                     devices_range Devs, const property_list &PropList,
+ProgramManager::link(device_images_range Imgs, devices_range Devs,
+                     const property_list &PropList,
                      bool AllowUnresolvedSymbols) {
   {
     auto NoAllowedPropertiesCheck = [](int) { return false; };
@@ -2989,17 +2987,13 @@ ProgramManager::link(sycl::span<const device_image_plain> Imgs,
         PropList, NoAllowedPropertiesCheck, NoAllowedPropertiesCheck);
   }
 
-  std::vector<ur_program_handle_t> URPrograms;
-  URPrograms.reserve(Imgs.size());
-  for (const device_image_plain &Img : Imgs)
-    URPrograms.push_back(getSyclObjImpl(Img)->get_ur_program());
-
+  auto URPrograms = Imgs.to<std::vector<ur_program_handle_t>>();
   auto URDevices = Devs.to<std::vector<ur_device_handle_t>>();
 
   // FIXME: Linker options are picked from the first object, but is that safe?
   std::string LinkOptionsStr;
   applyLinkOptionsFromEnvironment(LinkOptionsStr);
-  device_image_impl &FirstImgImpl = *getSyclObjImpl(Imgs[0]);
+  device_image_impl &FirstImgImpl = Imgs.front();
   if (LinkOptionsStr.empty() && FirstImgImpl.get_bin_image_ref())
     appendLinkOptionsFromImage(LinkOptionsStr,
                                *(FirstImgImpl.get_bin_image_ref()));
@@ -3061,8 +3055,8 @@ ProgramManager::link(sycl::span<const device_image_plain> Imgs,
     // underlying program disposed of). Protecting from incorrect values by
     // removal of map entries with same handle (obviously invalid entries).
     std::ignore = NativePrograms.erase(LinkedProg);
-    for (const device_image_plain &Img : Imgs) {
-      if (auto BinImageRef = getSyclObjImpl(Img)->get_bin_image_ref())
+    for (device_image_impl &Img : Imgs) {
+      if (auto BinImageRef = Img.get_bin_image_ref())
         NativePrograms.insert(
             {LinkedProg, {ContextImpl.shared_from_this(), BinImageRef}});
     }
@@ -3078,15 +3072,14 @@ ProgramManager::link(sycl::span<const device_image_plain> Imgs,
   KernelNameSetT MergedKernelNames;
   std::map<std::string, KernelArgMask, std::less<>>
       MergedEliminatedKernelArgMasks;
-  for (const device_image_plain &DevImg : Imgs) {
-    device_image_impl &DevImgImpl = *getSyclObjImpl(DevImg);
-    CombinedOrigins |= DevImgImpl.getOriginMask();
-    RTCInfoPtrs.emplace_back(&(DevImgImpl.getRTCInfo()));
-    MergedKernelNames.insert(DevImgImpl.getKernelNames().begin(),
-                             DevImgImpl.getKernelNames().end());
+  for (device_image_impl &DevImg : Imgs) {
+    CombinedOrigins |= DevImg.getOriginMask();
+    RTCInfoPtrs.emplace_back(&(DevImg.getRTCInfo()));
+    MergedKernelNames.insert(DevImg.getKernelNames().begin(),
+                             DevImg.getKernelNames().end());
     MergedEliminatedKernelArgMasks.insert(
-        DevImgImpl.getEliminatedKernelArgMasks().begin(),
-        DevImgImpl.getEliminatedKernelArgMasks().end());
+        DevImg.getEliminatedKernelArgMasks().begin(),
+        DevImg.getEliminatedKernelArgMasks().end());
   }
   auto MergedRTCInfo = detail::KernelCompilerBinaryInfo::Merge(RTCInfoPtrs);
 
@@ -3100,16 +3093,12 @@ ProgramManager::link(sycl::span<const device_image_plain> Imgs,
       std::move(MergedImageStorage)))};
 }
 
-void ProgramManager::dynamicLink(sycl::span<const device_image_plain> Imgs) {
+void ProgramManager::dynamicLink(device_images_range Imgs) {
   if (Imgs.empty())
     return;
 
-  std::vector<ur_program_handle_t> URPrograms;
-  URPrograms.reserve(Imgs.size());
-  for (const device_image_plain &Img : Imgs)
-    URPrograms.push_back(getSyclObjImpl(Img)->get_ur_program());
-
-  device_image_impl &FirstImgImpl = *getSyclObjImpl(Imgs[0]);
+  auto URPrograms = Imgs.to<std::vector<ur_program_handle_t>>();
+  device_image_impl &FirstImgImpl = Imgs.front();
   auto [URCtx, Adapter] =
       get_ur_handles(*getSyclObjImpl(FirstImgImpl.get_context()));
   Adapter->call<UrApiKind::urProgramDynamicLinkExp>(URCtx, URPrograms.size(),
diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp
index 3d2f5a8974349..ee10938f1b418 100644
--- a/sycl/source/detail/program_manager/program_manager.hpp
+++ b/sycl/source/detail/program_manager/program_manager.hpp
@@ -25,7 +25,6 @@
 #include <sycl/detail/util.hpp>
 #include <sycl/device.hpp>
 #include <sycl/kernel_bundle.hpp>
-#include <sycl/sycl_span.hpp>
 
 #include <array>
 #include <cstdint>
@@ -77,6 +76,7 @@ class device_impl;
 class devices_range;
 class queue_impl;
 class event_impl;
+class device_images_range;
 // DeviceLibExt is shared between sycl runtime and sycl-post-link tool.
 // If any update is made here, need to sync with DeviceLibExt definition
 // in llvm/tools/sycl-post-link/sycl-post-link.cpp
@@ -356,12 +356,13 @@ class ProgramManager {
 
   // Produces set of device images by convering input device images to object
   // the executable state
-  std::vector<device_image_plain>
-  link(sycl::span<const device_image_plain> Imgs, devices_range Devs,
-       const property_list &PropList, bool AllowUnresolvedSymbols = false);
+  std::vector<device_image_plain> link(device_images_range Imgs,
+                                       devices_range Devs,
+                                       const property_list &PropList,
+                                       bool AllowUnresolvedSymbols = false);
 
   // Dynamically links images in executable state.
-  void dynamicLink(sycl::span<const device_image_plain> Imgs);
+  void dynamicLink(device_images_range Imgs);
 
   // Produces new device image by converting input device image to the
   // executable state