diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 1bbfb4d96f46d..5787bc6c9af59 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -495,7 +495,21 @@ class kernel_bundle_impl { std::vector DeviceVec; DeviceVec.reserve(Devices.size()); for (const auto &SyclDev : Devices) { - ur_device_handle_t Dev = getSyclObjImpl(SyclDev)->getHandleRef(); + DeviceImplPtr DevImpl = getSyclObjImpl(SyclDev); + if (!ContextImpl->hasDevice(DevImpl)) { + throw sycl::exception(make_error_code(errc::invalid), + "device not part of kernel_bundle context"); + } + if (!DevImpl->extOneapiCanCompile(MLanguage)) { + // This error cannot not be exercised in the current implementation, as + // compatibility with a source language depends on the backend's + // capabilities and all devices in one context share the same backend in + // the current implementation, so this would lead to an error already + // during construction of the source bundle. + throw sycl::exception(make_error_code(errc::invalid), + "device does not support source language"); + } + ur_device_handle_t Dev = DevImpl->getHandleRef(); DeviceVec.push_back(Dev); } diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_basic.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_basic.cpp index a2352977f0c2c..ce6b251b9080a 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_basic.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_basic.cpp @@ -1,4 +1,4 @@ -//==- kernel_compiler_sycl_jit.cpp --- kernel_compiler extension tests -----==// +//==---- kernel_compiler_basic.cpp --- kernel_compiler extension tests -----==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_context_error.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_context_error.cpp new file mode 100644 index 0000000000000..e33d85365924e --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_context_error.cpp @@ -0,0 +1,85 @@ +//==- kernel_compiler_context_error.cpp -- kernel_compiler extension tests -==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// REQUIRES: (opencl || level_zero) + +// UNSUPPORTED: accelerator +// UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there. + +// RUN: %{build} -o %t.out +// RUN: %{run-unfiltered-devices} %t.out + +#include +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +const std::string source = R"""( + #include + namespace syclext = sycl::ext::oneapi; + namespace syclexp = sycl::ext::oneapi::experimental; + + extern "C" + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) + void iota(float start, float *ptr) { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + ptr[id] = start + static_cast(id); + } +)"""; + +int main() { + + auto has_multiple_compatible_devices = [](sycl::platform platform) -> bool { + auto devices = platform.get_devices(); + if (devices.size() < 2) { + return false; + } + for (auto dev : devices) { + if (!dev.ext_oneapi_can_compile(syclexp::source_language::sycl)) { + return false; + } + } + return true; + }; + + std::vector all_devices = [&]() -> std::vector { + for (auto platform : sycl::platform::get_platforms()) { + if (has_multiple_compatible_devices(platform)) { + return platform.get_devices(); + } + } + return {}; + }(); + + if (all_devices.size() < 2) { + std::cerr << "Cannot find platform with more than 1 device, skipping" + << std::endl; + return 0; + } + + sycl::context single_device_context{all_devices.front()}; + + // Create a source kernel bundle with a context that contains only one device. + sycl::kernel_bundle kb_src = + syclexp::create_kernel_bundle_from_source( + single_device_context, syclexp::source_language::sycl, source); + + // Compile the kernel. There is no need to use the "registered_names" + // property because the kernel is declared extern "C". + try { + syclexp::build(kb_src, all_devices); + assert(false && "out-of-context device not detected"); + } catch (sycl::exception &e) { + assert(e.code() == sycl::errc::invalid); + assert(std::string(e.what()).find( + "device not part of kernel_bundle context") != + std::string::npos); + } + return 0; +} diff --git a/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp b/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp index 23ea51be5d7c4..cecce9368f21b 100644 --- a/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp +++ b/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp @@ -6,7 +6,7 @@ // CHECK-DAG: README.md // CHECK-DAG: lit.cfg.py // -// CHECK-NUM-MATCHES: 10 +// CHECK-NUM-MATCHES: 11 // // This test verifies that `` isn't used in E2E tests. Instead, // fine-grained includes should used, see