[SYCL] Fix for native_specialization_constant() (#9085)

If MBinImage is nullptr, it leads to SIGSEGV.
This patch fixes ProgramManager::link() function, so MBinImage will not
be nullptr.
Also changes link() signature to accept a single device image instead of
vector,
because we use it to pass the vector of one image, so it makes no sense.

---------

Co-authored-by: Steffen Larsen <[email protected]>

Author: KornevNikita

sycl/source/detail/device_image_impl.hpp

@@ -112,6 +112,15 @@ class device_image_impl {
   bool all_specialization_constant_native() const noexcept {
     // Specialization constants are natively supported in JIT mode on backends,
     // that are using SPIR-V as IR
+
+    // Not sure if it's possible currently, but probably it may happen if the
+    // kernel bundle is created with interop function. Now the only one such
+    // function is make_kernel(), but I'm not sure if it's even possible to
+    // use spec constant with such kernel. So, in such case we need to check
+    // if it's JIT or no somehow.
+    assert(MBinImage &&
+           "native_specialization_constant() called for unimplemented case");
+
     auto IsJITSPIRVTarget = [](const char *Target) {
       return (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64) == 0 ||
               strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV32) == 0);

sycl/source/detail/kernel_bundle_impl.hpp

@@ -213,10 +213,9 @@ class kernel_bundle_impl {
                          }))
           continue;
 
-        const std::vector<device_image_plain> VectorOfOneImage{DeviceImage};
         std::vector<device_image_plain> LinkedResults =
-            detail::ProgramManager::getInstance().link(VectorOfOneImage,
-                                                       MDevices, PropList);
+            detail::ProgramManager::getInstance().link(DeviceImage, MDevices,
+                                                       PropList);
         MDeviceImages.insert(MDeviceImages.end(), LinkedResults.begin(),
                              LinkedResults.end());
       }

sycl/source/detail/program_manager/program_manager.cpp

@@ -2105,15 +2105,13 @@ ProgramManager::compile(const device_image_plain &DeviceImage,
 }
 
 std::vector<device_image_plain>
-ProgramManager::link(const std::vector<device_image_plain> &DeviceImages,
+ProgramManager::link(const device_image_plain &DeviceImage,
                      const std::vector<device> &Devs,
                      const property_list &PropList) {
   (void)PropList;
 
   std::vector<pi_program> PIPrograms;
-  PIPrograms.reserve(DeviceImages.size());
-  for (const device_image_plain &DeviceImage : DeviceImages)
-    PIPrograms.push_back(getSyclObjImpl(DeviceImage)->get_program_ref());
+  PIPrograms.push_back(getSyclObjImpl(DeviceImage)->get_program_ref());
 
   std::vector<pi_device> PIDevices;
   PIDevices.reserve(Devs.size());
@@ -2123,14 +2121,12 @@ ProgramManager::link(const std::vector<device_image_plain> &DeviceImages,
   std::string LinkOptionsStr;
   applyLinkOptionsFromEnvironment(LinkOptionsStr);
   if (LinkOptionsStr.empty()) {
-    for (const device_image_plain &DeviceImage : DeviceImages) {
-      const std::shared_ptr<device_image_impl> &InputImpl =
-          getSyclObjImpl(DeviceImage);
-      appendLinkOptionsFromImage(LinkOptionsStr,
-                                 *(InputImpl->get_bin_image_ref()));
-    }
+    const std::shared_ptr<device_image_impl> &InputImpl =
+        getSyclObjImpl(DeviceImage);
+    appendLinkOptionsFromImage(LinkOptionsStr,
+                               *(InputImpl->get_bin_image_ref()));
   }
-  const context &Context = getSyclObjImpl(DeviceImages[0])->get_context();
+  const context &Context = getSyclObjImpl(DeviceImage)->get_context();
   const ContextImplPtr ContextImpl = getSyclObjImpl(Context);
   const detail::plugin &Plugin = ContextImpl->getPlugin();
 
@@ -2152,55 +2148,53 @@ ProgramManager::link(const std::vector<device_image_plain> &DeviceImages,
   std::shared_ptr<std::vector<kernel_id>> KernelIDs{new std::vector<kernel_id>};
   std::vector<unsigned char> NewSpecConstBlob;
   device_image_impl::SpecConstMapT NewSpecConstMap;
-  for (const device_image_plain &DeviceImage : DeviceImages) {
-    std::shared_ptr<device_image_impl> DeviceImageImpl =
-        getSyclObjImpl(DeviceImage);
 
-    // Duplicates are not expected here, otherwise piProgramLink should fail
-    KernelIDs->insert(KernelIDs->end(),
-                      DeviceImageImpl->get_kernel_ids_ptr()->begin(),
-                      DeviceImageImpl->get_kernel_ids_ptr()->end());
+  std::shared_ptr<device_image_impl> DeviceImageImpl =
+      getSyclObjImpl(DeviceImage);
 
-    // 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());
-
-      // 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()) {
-        std::vector<device_image_impl::SpecConstDescT> &NewDescEntries =
-            NewSpecConstMap[SpecConstIt.first];
-        assert(NewDescEntries.empty() &&
-               "Specialization constant already exists in the map.");
-        NewDescEntries.reserve(SpecConstIt.second.size());
-        for (const device_image_impl::SpecConstDescT &SpecConstDesc :
-             SpecConstIt.second) {
-          device_image_impl::SpecConstDescT NewSpecConstDesc = SpecConstDesc;
-          NewSpecConstDesc.BlobOffset += NewSpecConstBlob.size();
-          NewDescEntries.push_back(std::move(NewSpecConstDesc));
-        }
-      }
+  // Duplicates are not expected here, otherwise piProgramLink should fail
+  KernelIDs->insert(KernelIDs->end(),
+                    DeviceImageImpl->get_kernel_ids_ptr()->begin(),
+                    DeviceImageImpl->get_kernel_ids_ptr()->end());
 
-      // 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());
+  // 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());
+
+    // 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()) {
+      std::vector<device_image_impl::SpecConstDescT> &NewDescEntries =
+          NewSpecConstMap[SpecConstIt.first];
+      assert(NewDescEntries.empty() &&
+             "Specialization constant already exists in the map.");
+      NewDescEntries.reserve(SpecConstIt.second.size());
+      for (const device_image_impl::SpecConstDescT &SpecConstDesc :
+           SpecConstIt.second) {
+        device_image_impl::SpecConstDescT NewSpecConstDesc = SpecConstDesc;
+        NewSpecConstDesc.BlobOffset += NewSpecConstBlob.size();
+        NewDescEntries.push_back(std::move(NewSpecConstDesc));
+      }
     }
+
+    // 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());
   }
+
   // device_image_impl expects kernel ids to be sorted for fast search
   std::sort(KernelIDs->begin(), KernelIDs->end(), LessByHash<kernel_id>{});
 
+  auto BinImg = getSyclObjImpl(DeviceImage)->get_bin_image_ref();
   DeviceImageImplPtr ExecutableImpl =
       std::make_shared<detail::device_image_impl>(
-          /*BinImage=*/nullptr, Context, Devs, bundle_state::executable,
-          std::move(KernelIDs), LinkedProg, std::move(NewSpecConstMap),
-          std::move(NewSpecConstBlob));
+          BinImg, Context, Devs, bundle_state::executable, std::move(KernelIDs),
+          LinkedProg, std::move(NewSpecConstMap), std::move(NewSpecConstBlob));
 
   // TODO: Make multiple sets of device images organized by devices they are
   // compiled for.

sycl/source/detail/program_manager/program_manager.hpp

@@ -279,9 +279,9 @@ class ProgramManager {
 
   // Produces set of device images by convering input device images to object
   // the executable state
-  std::vector<device_image_plain>
-  link(const std::vector<device_image_plain> &DeviceImages,
-       const std::vector<device> &Devs, const property_list &PropList);
+  std::vector<device_image_plain> link(const device_image_plain &DeviceImages,
+                                       const std::vector<device> &Devs,
+                                       const property_list &PropList);
 
   // Produces new device image by converting input device image to the
   // executable state

sycl/test-e2e/SpecConstants/2020/kernel-bundle-api.cpp

@@ -266,9 +266,32 @@ bool test_set_and_get_on_device(sycl::queue q) {
 }
 
 bool test_native_specialization_constant(sycl::queue q) {
+  {
+    q.submit([&](sycl::handler &cgh) {
+      cgh.single_task<class Kernel>([=](sycl::kernel_handler h) {
+        h.get_specialization_constant<int_id>();
+      });
+    });
+
+    auto inputBundle =
+        sycl::get_kernel_bundle<class Kernel, sycl::bundle_state::input>(
+            q.get_context(), {q.get_device()});
+    auto objectBundle = sycl::compile(inputBundle);
+    auto execBundleViaLink = sycl::link(objectBundle);
+    auto BE = q.get_backend();
+    bool expected = (BE == sycl::backend::opencl ||
+                     BE == sycl::backend::ext_oneapi_level_zero)
+                        ? true
+                        : false;
+    if (!check_value(expected,
+                     execBundleViaLink.native_specialization_constant(),
+                     "linked bundle native specialization constant"))
+      return false;
+  }
+
   const auto always_false_selector = [](auto device_image) { return false; };
   auto bundle = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
       q.get_context(), always_false_selector);
-  return check_value(bundle.native_specialization_constant(), false,
+  return check_value(false, bundle.native_specialization_constant(),
                      "empty bundle native specialization constant");
 }