[SYCL][Docs] Disallow sycl::join with source-based kernel_bundle (#18282)

This commit adds a constraint to sycl::join specifying that it is not
available when the kernel bundle state is source-based.

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc

@@ -642,6 +642,14 @@ As a result, the only `kernel_bundle` member functions from the core SYCL
 specification that are available for bundles in `ext_oneapi_source` state are
 `get_backend`, `get_context`, and `get_devices`.
 
+=== New constraint for `join` function
+
+This extension adds the following constraint to the `join` functions from the
+core SYCL specification:
+
+> _Constraints:_ This function is not available when `State` is
+> `bundle_state::ext_oneapi_source`.
+
 === Interaction with existing kernel bundle member functions
 
 Kernels created from online compilation of source code do not have any

sycl/include/sycl/kernel_bundle.hpp

@@ -885,7 +885,8 @@ join_impl(const std::vector<detail::KernelBundleImplPtr> &Bundles,
 /// \returns a new kernel bundle that represents the union of all the device
 /// images in the input bundles with duplicates removed.
 template <sycl::bundle_state State>
-sycl::kernel_bundle<State>
+std::enable_if_t<State != sycl::bundle_state::ext_oneapi_source,
+                 sycl::kernel_bundle<State>>
 join(const std::vector<sycl::kernel_bundle<State>> &Bundles) {
   // Convert kernel_bundle<State> to impls to abstract template parameter away
   std::vector<detail::KernelBundleImplPtr> KernelBundleImpls;

sycl/test-e2e/KernelCompiler/sycl_join.cpp

@@ -81,46 +81,6 @@ int main() {
   source_kb KBSrc2 = syclex::create_kernel_bundle_from_source(
       Ctx, syclex::source_language::sycl, SYCLSource2);
 
-  // Test joining of source kernel bundles.
-  {
-    std::vector<source_kb> KBSrcs{KBSrc1, KBSrc2};
-    source_kb KBSrcJoined = sycl::join(KBSrcs);
-
-    exe_kb KBExeJoined = syclex::build(KBSrcJoined);
-    assert(KBExeJoined.ext_oneapi_has_kernel("TestKernel1"));
-    assert(KBExeJoined.ext_oneapi_has_kernel("TestKernel2"));
-
-    sycl::kernel K1 = KBExeJoined.ext_oneapi_get_kernel("TestKernel1");
-    sycl::kernel K2 = KBExeJoined.ext_oneapi_get_kernel("TestKernel2");
-
-    int *IntPtr = sycl::malloc_shared<int>(1, Q);
-    *IntPtr = 0;
-
-    Q.submit([&](sycl::handler &CGH) {
-       CGH.set_args(IntPtr);
-       CGH.single_task(K1);
-     }).wait_and_throw();
-
-    if (*IntPtr != 42) {
-      std::cout << "TestKernel1 in joined source bundles failed: " << *IntPtr
-                << " != 42\n";
-      ++Failed;
-    }
-
-    Q.submit([&](sycl::handler &CGH) {
-       CGH.set_args(IntPtr);
-       CGH.single_task(K2);
-     }).wait_and_throw();
-
-    if (*IntPtr != 24) {
-      std::cout << "TestKernel1 in joined source bundles failed: " << *IntPtr
-                << " != 24\n";
-      ++Failed;
-    }
-
-    sycl::free(IntPtr, Q);
-  }
-
   auto KBExe1 = std::make_shared<exe_kb>(syclex::build(KBSrc1));
   auto KBExe2 = std::make_shared<exe_kb>(syclex::build(KBSrc2));