diff --git a/SYCL/Regression/kernel_bundle_ignore_sycl_external.cpp b/SYCL/Regression/kernel_bundle_ignore_sycl_external.cpp new file mode 100644 index 0000000000..44cf661a05 --- /dev/null +++ b/SYCL/Regression/kernel_bundle_ignore_sycl_external.cpp @@ -0,0 +1,40 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out +// +// XFAIL: cuda || hip + +#include + +class KernelName; + +SYCL_EXTERNAL +int f(int a) { return a + 1; } + +int main() { + const sycl::device Dev{sycl::default_selector{}}; + const sycl::context Ctx{Dev}; + sycl::queue Q{Ctx, Dev}; + + assert(sycl::get_kernel_ids().size() == 1); + + sycl::kernel_bundle EmptyKernelBundle = + sycl::get_kernel_bundle(Ctx, {Dev}, {}); + + assert(EmptyKernelBundle.get_kernel_ids().size() == 0); + + sycl::kernel_bundle KernelBundle = + sycl::get_kernel_bundle(Ctx, {Dev}); + sycl::kernel_id KernelID = sycl::get_kernel_id(); + + assert(KernelBundle.get_kernel_ids().size() == 1); + assert(KernelBundle.has_kernel(KernelID)); + + cl::sycl::buffer Buf(sycl::range<1>{1}); + Q.submit([&](sycl::handler &CGH) { + auto Acc = Buf.get_access(CGH); + CGH.use_kernel_bundle(KernelBundle); + CGH.single_task([=]() { Acc[0] = 42; }); + }); +}