From b448e8ac809d0ca0f21bd7584eca23773166ede2 Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Thu, 6 Mar 2025 13:26:28 +0000 Subject: [PATCH 1/2] [SYCL][RTC] Implement specified error behavior Detect error scenarios in `sycl::ext::oneapi::build` and throw the errors specified in the extension. This handles two scenarios: - Calling build where one of the devices given in the list does not support the source language of the bundle. - Calling build where one of the devices given in the list does not belong to the context of the source bundle. Signed-off-by: Lukas Sommer --- sycl/source/detail/kernel_bundle_impl.hpp | 16 +++- .../KernelCompiler/kernel_compiler_basic.cpp | 2 +- .../kernel_compiler_context_error.cpp | 85 +++++++++++++++++++ 3 files changed, 101 insertions(+), 2 deletions(-) create mode 100644 sycl/test-e2e/KernelCompiler/kernel_compiler_context_error.cpp 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..6fe0c7c0a473f --- /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_jit, 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; +} From a5362da3347bf47cc0956b3075bf01241689cad7 Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Fri, 7 Mar 2025 08:21:05 +0000 Subject: [PATCH 2/2] Address PR feedback Signed-off-by: Lukas Sommer --- sycl/test-e2e/KernelCompiler/kernel_compiler_context_error.cpp | 2 +- sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_context_error.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_context_error.cpp index 6fe0c7c0a473f..e33d85365924e 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_context_error.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_context_error.cpp @@ -68,7 +68,7 @@ int main() { // 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_jit, 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". 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