[NFCI][SYCL] Simplify retrieveKernelBinary

Not sure if side effects of `createURProgram` were necessary and even if
they were, don't know if I'm removing them or just delaying till the
next access of some cache.

Author: aelovikov-intel

sycl/source/detail/helpers.cpp

@@ -36,9 +36,9 @@ markBufferAsInternal(const std::shared_ptr<buffer_impl> &BufImpl) {
   BufImpl->markAsInternal();
 }
 
-std::tuple<const RTDeviceBinaryImage *, ur_program_handle_t>
-retrieveKernelBinary(queue_impl &Queue, KernelNameStrRefT KernelName,
-                     CGExecKernel *KernelCG) {
+const RTDeviceBinaryImage *retrieveKernelBinary(queue_impl &Queue,
+                                                KernelNameStrRefT KernelName,
+                                                CGExecKernel *KernelCG) {
   device_impl &Dev = Queue.getDeviceImpl();
   bool isNvidia = Dev.getBackend() == backend::ext_oneapi_cuda;
   bool isHIP = Dev.getBackend() == backend::ext_oneapi_hip;
@@ -55,36 +55,22 @@ retrieveKernelBinary(queue_impl &Queue, KernelNameStrRefT KernelName,
           return DI->getFormat() == SYCL_DEVICE_BINARY_TYPE_LLVMIR_BITCODE &&
                  DI->getRawData().DeviceTargetSpec == TargetSpec;
         });
-    if (DeviceImage == DeviceImages.end()) {
-      return {nullptr, nullptr};
-    }
-    context_impl &ContextImpl = Queue.getContextImpl();
-    ur_program_handle_t Program =
-        detail::ProgramManager::getInstance().createURProgram(
-            **DeviceImage, ContextImpl, {createSyclObjFromImpl<device>(Dev)});
-    return {*DeviceImage, Program};
+    assert(DeviceImage != DeviceImages.end() &&
+           "Failed to obtain a binary image.");
+    return *DeviceImage;
   }
 
-  const RTDeviceBinaryImage *DeviceImage = nullptr;
-  ur_program_handle_t Program = nullptr;
-  auto KernelBundleImpl = KernelCG->getKernelBundle();
-  if (KernelCG->MSyclKernel != nullptr) {
-    DeviceImage = KernelCG->MSyclKernel->getDeviceImage()->get_bin_image_ref();
-    Program = KernelCG->MSyclKernel->getDeviceImage()->get_ur_program_ref();
-  } else if (auto SyclKernelImpl =
-                 KernelBundleImpl ? KernelBundleImpl->tryGetKernel(KernelName)
-                                  : std::shared_ptr<kernel_impl>{nullptr}) {
-    // Retrieve the device image from the kernel bundle.
-    DeviceImage = SyclKernelImpl->getDeviceImage()->get_bin_image_ref();
-    Program = SyclKernelImpl->getDeviceImage()->get_ur_program_ref();
-  } else {
-    context_impl &ContextImpl = Queue.getContextImpl();
-    DeviceImage = &detail::ProgramManager::getInstance().getDeviceImage(
-        KernelName, ContextImpl, Dev);
-    Program = detail::ProgramManager::getInstance().createURProgram(
-        *DeviceImage, ContextImpl, {createSyclObjFromImpl<device>(Dev)});
-  }
-  return {DeviceImage, Program};
+  if (KernelCG->MSyclKernel != nullptr)
+    return KernelCG->MSyclKernel->getDeviceImage()->get_bin_image_ref();
+
+  if (auto KernelBundleImpl = KernelCG->getKernelBundle())
+    if (KernelBundleImpl->tryGetKernel(KernelName))
+      // Retrieve the device image from the kernel bundle.
+      return SyclKernelImpl->getDeviceImage()->get_bin_image_ref();
+
+  context_impl &ContextImpl = Queue.getContextImpl();
+  return &detail::ProgramManager::getInstance().getDeviceImage(
+      KernelName, ContextImpl, Dev);
 }
 
 } // namespace detail

sycl/source/detail/helpers.hpp

@@ -27,7 +27,7 @@ class RTDeviceBinaryImage;
 void waitEvents(std::vector<sycl::event> DepEvents);
 #endif
 
-std::tuple<const RTDeviceBinaryImage *, ur_program_handle_t>
+const RTDeviceBinaryImage *
 retrieveKernelBinary(queue_impl &Queue, KernelNameStrRefT KernelName,
                      CGExecKernel *CGKernel = nullptr);
 } // namespace detail

sycl/source/detail/scheduler/commands.cpp

@@ -3258,8 +3258,7 @@ ur_result_t ExecCGCommand::enqueueImpQueue() {
 
     const RTDeviceBinaryImage *BinImage = nullptr;
     if (detail::SYCLConfig<detail::SYCL_JIT_AMDGCN_PTX_KERNELS>::get()) {
-      std::tie(BinImage, std::ignore) =
-          retrieveKernelBinary(*MQueue, KernelName);
+      BinImage = retrieveKernelBinary(*MQueue, KernelName);
       assert(BinImage && "Failed to obtain a binary image.");
     }
     enqueueImpKernel(

sycl/source/handler.cpp

@@ -597,7 +597,7 @@ event handler::finalize() {
 #endif
         const detail::RTDeviceBinaryImage *BinImage = nullptr;
         if (detail::SYCLConfig<detail::SYCL_JIT_AMDGCN_PTX_KERNELS>::get()) {
-          std::tie(BinImage, std::ignore) = detail::retrieveKernelBinary(
+          BinImage = detail::retrieveKernelBinary(
               impl->get_queue(), toKernelNameStrT(MKernelName));
           assert(BinImage && "Failed to obtain a binary image.");
         }