diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index b8970ea8cb1c0..9328054445fcb 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -246,25 +246,20 @@ class kernel_bundle_impl // Due to a bug in L0, specializations with conflicting IDs will overwrite // each other when linked together, so to avoid this issue we link - // images with specialization constants in separation. + // regular offline-compiled SYCL device images in separation. // TODO: Remove when spec const overwriting issue has been fixed in L0. - std::vector ImagesWithSpecConsts; - std::unordered_set ImagesWithSpecConstsSet; + std::vector OfflineDeviceImages; + std::unordered_set OfflineDeviceImageSet; for (const kernel_bundle &ObjectBundle : ObjectBundles) { for (const DevImgPlainWithDeps &DeviceImageWithDeps : getSyclObjImpl(ObjectBundle)->MDeviceImages) { - if (std::none_of(DeviceImageWithDeps.begin(), DeviceImageWithDeps.end(), - [](const device_image_plain &DevImg) { - const RTDeviceBinaryImage *BinImg = - getSyclObjImpl(DevImg)->get_bin_image_ref(); - return BinImg && BinImg->getSpecConstants().size(); - })) - continue; - - ImagesWithSpecConsts.push_back(&DeviceImageWithDeps); - for (const device_image_plain &DevImg : DeviceImageWithDeps) - ImagesWithSpecConstsSet.insert(&*getSyclObjImpl(DevImg)); + if (getSyclObjImpl(DeviceImageWithDeps.getMain())->getOriginMask() & + ImageOriginSYCLOffline) { + OfflineDeviceImages.push_back(&DeviceImageWithDeps); + for (const device_image_plain &DevImg : DeviceImageWithDeps) + OfflineDeviceImageSet.insert(&*getSyclObjImpl(DevImg)); + } } } @@ -274,22 +269,16 @@ class kernel_bundle_impl std::set DevImagesSet; std::unordered_set SeenBinImgs; for (const kernel_bundle &ObjectBundle : - ObjectBundles) { + ObjectBundles) for (device_image_impl &DevImg : - getSyclObjImpl(ObjectBundle)->device_images()) { - const RTDeviceBinaryImage *BinImg = DevImg.get_bin_image_ref(); - // We have duplicate images if either the underlying binary image has - // been seen before or the device image implementation is in the - // image set already. - if ((BinImg && SeenBinImgs.count(BinImg)) || - ImagesWithSpecConstsSet.count(&DevImg)) - continue; - SeenBinImgs.insert(BinImg); - DevImagesSet.insert(&DevImg); - } - } - DevImages = device_images_range{DevImagesSet} - .to>(); + getSyclObjImpl(ObjectBundle)->device_images()) + if (OfflineDeviceImageSet.find(&DevImg) == + OfflineDeviceImageSet.end()) + DevImagesSet.insert(&DevImg); + DevImages.reserve(DevImagesSet.size()); + for (auto It = DevImagesSet.begin(); It != DevImagesSet.end();) + DevImages.push_back(createSyclObjFromImpl( + *DevImagesSet.extract(It++).value())); } // Check for conflicting kernels in RTC kernel bundles. @@ -392,8 +381,7 @@ class kernel_bundle_impl } // ... And link the offline images in separation. (Workaround.) - for (const DevImgPlainWithDeps *DeviceImageWithDeps : - ImagesWithSpecConsts) { + for (const DevImgPlainWithDeps *DeviceImageWithDeps : OfflineDeviceImages) { // Skip images which are not compatible with devices provided if (std::none_of(get_devices().begin(), get_devices().end(), [DeviceImageWithDeps](device_impl &Dev) { diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index d5e9aa0b3e5a6..3346b932260cf 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -2587,10 +2587,7 @@ ProgramManager::getSYCLDeviceImagesWithCompatibleState( std::lock_guard KernelIDsGuard(m_KernelIDsMutex); ImgInfo.KernelIDs = m_BinImg2KernelIDs[BinImage]; } - ImgInfo.Deps = - collectDeviceImageDeps(*BinImage, Dev, - /*ErrorOnUnresolvableImport=*/TargetState == - bundle_state::executable); + ImgInfo.Deps = collectDeviceImageDeps(*BinImage, Dev); } const bundle_state ImgState = ImgInfo.State; const std::shared_ptr> &ImageKernelIDs = diff --git a/sycl/test-e2e/SYCLBIN/link_mixed_opt_input.cpp b/sycl/test-e2e/SYCLBIN/link_mixed_opt_input.cpp index 06d220dbe20f4..072b8ac903180 100644 --- a/sycl/test-e2e/SYCLBIN/link_mixed_opt_input.cpp +++ b/sycl/test-e2e/SYCLBIN/link_mixed_opt_input.cpp @@ -20,7 +20,7 @@ // RUN: %clangxx --offload-new-driver -fsyclbin=input -fsycl-allow-device-image-dependencies -O0 %S/Inputs/exporting_function.cpp -o %t.export.syclbin // RUN: %clangxx --offload-new-driver -fsyclbin=input -fsycl-allow-device-image-dependencies -O1 %S/Inputs/importing_kernel.cpp -o %t.import.syclbin // RUN: %{build} -o %t.out -// RUN: %{l0_leak_check} %{run} %t.out %t.export.syclbin %t.import.syclbin +// RUN: %{run} %t.out %t.export.syclbin %t.import.syclbin #define SYCLBIN_INPUT_STATE diff --git a/sycl/test-e2e/SYCLBIN/link_mixed_opt_object.cpp b/sycl/test-e2e/SYCLBIN/link_mixed_opt_object.cpp index 66c52a6192e2a..4cf5e197c98a4 100644 --- a/sycl/test-e2e/SYCLBIN/link_mixed_opt_object.cpp +++ b/sycl/test-e2e/SYCLBIN/link_mixed_opt_object.cpp @@ -20,7 +20,7 @@ // RUN: %clangxx --offload-new-driver -fsyclbin=object -fsycl-allow-device-image-dependencies -O0 %S/Inputs/exporting_function.cpp -o %t.export.syclbin // RUN: %clangxx --offload-new-driver -fsyclbin=object -fsycl-allow-device-image-dependencies -O1 %S/Inputs/importing_kernel.cpp -o %t.import.syclbin // RUN: %{build} -o %t.out -// RUN: %{l0_leak_check} %{run} %t.out %t.export.syclbin %t.import.syclbin +// RUN: %{run} %t.out %t.export.syclbin %t.import.syclbin #define SYCLBIN_OBJECT_STATE diff --git a/sycl/unittests/kernel-and-program/OutOfResources.cpp b/sycl/unittests/kernel-and-program/OutOfResources.cpp index e55645eb5a315..75534205a4b9b 100644 --- a/sycl/unittests/kernel-and-program/OutOfResources.cpp +++ b/sycl/unittests/kernel-and-program/OutOfResources.cpp @@ -173,7 +173,7 @@ TEST_P(OutOfResourcesTestSuite, urProgramLink) { auto b3 = sycl::link({b1, b2}); EXPECT_FALSE(outOfResourcesToggle); // one restart due to out of resources, one link per each of b1 and b2. - EXPECT_EQ(nProgramLink, 2); + EXPECT_EQ(nProgramLink, 3); // no programs should be in the cache due to out of resources. { detail::KernelProgramCache::ProgramCache &Cache =