[SYCL] Fix linking of kernel-bundles (#4398)

when fsycl-device-code-split=per_kernel is used in conjunction with
multiple kernels, then explicit sycl::link(sycl::compile(my-kernel-bundle))
will fail, because we link all device images in a kernel together. But, the
device-images are independent, there is presently no known situation
where they should be linked together. So this PR works around this
limitation by linking them separately. This may have to be revisited once
shared library style linking is supported, but that is likely true whether
this bug is fixed or not.

Signed-off-by: Chris Perkins <[email protected]>

Author: cperkinsintel

sycl/source/detail/kernel_bundle_impl.hpp

@@ -192,7 +192,13 @@ class kernel_bundle_impl {
     // TODO: Unify with c'tor for sycl::comile and sycl::build by calling
     // sycl::join on vector of kernel_bundles
 
-    std::vector<device_image_plain> DeviceImages;
+    // The loop below just links each device image separately, not linking any
+    // two device images together. This is correct so long as each device image
+    // has no unresolved symbols. That's the case when device images are created
+    // from generic SYCL APIs. There's no way in generic SYCL to create a kernel
+    // which references an undefined symbol. If we decide in the future to allow
+    // a backend interop API to create a "sycl::kernel_bundle" that references
+    // undefined symbols, then the logic in this loop will need to be changed.
     for (const kernel_bundle<bundle_state::object> &ObjectBundle :
          ObjectBundles) {
       for (const device_image_plain &DeviceImage : ObjectBundle) {
@@ -205,13 +211,15 @@ class kernel_bundle_impl {
                          }))
           continue;
 
-        DeviceImages.insert(DeviceImages.end(), DeviceImage);
+        const std::vector<device_image_plain> VectorOfOneImage{DeviceImage};
+        std::vector<device_image_plain> LinkedResults =
+            detail::ProgramManager::getInstance().link(VectorOfOneImage,
+                                                       MDevices, PropList);
+        MDeviceImages.insert(MDeviceImages.end(), LinkedResults.begin(),
+                             LinkedResults.end());
       }
     }
 
-    MDeviceImages = detail::ProgramManager::getInstance().link(
-        std::move(DeviceImages), MDevices, PropList);
-
     for (const kernel_bundle<bundle_state::object> &Bundle : ObjectBundles) {
       const KernelBundleImplPtr BundlePtr = getSyclObjImpl(Bundle);
       for (const std::pair<const std::string, std::vector<unsigned char>>