From fd45635c92114d9a219a0bb12ad293095061632f Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Wed, 4 Nov 2026 13:11:08 +0100 Subject: [PATCH 1/3] Pass the SYCL submission error to SyclKernelSubmitError DPCTLQueue_SubmitRange and DPCTLQueue_SubmitNDRange catch the exception thrown by sycl::queue::submit and return NULL, so Python only got "Kernel submission to Sycl queue failed." and the reason was printed only with DPCTL_VERBOSITY set. Both functions now keep the reason in a thread-local string, available through the new DPCTLQueue_GetLastSubmitError C-API function, and SyclQueue.submit_async appends it to the SyclKernelSubmitError message. Also document in create_kernel_bundle_from_sycl_source that, starting with DPC++ 2026.2, kernels launched over more than INT_MAX work-items must be compiled with -fno-sycl-id-queries-fit-in-int in copts. --- CHANGELOG.md | 2 + dpctl/_backend.pxd | 1 + dpctl/_sycl_queue.pyx | 14 ++- dpctl/compiler/_compiler.pyx | 10 +++ dpctl/tests/test_sycl_kernel_submit.py | 25 ++++++ .../dpctl_sycl_queue_interface.h | 17 ++++ .../source/dpctl_sycl_queue_interface.cpp | 90 +++++++++++-------- .../tests/test_sycl_queue_submit.cpp | 45 ++++++++++ 8 files changed, 163 insertions(+), 41 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 4b3d0c90d1..3b3c31e442 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -7,6 +7,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 ## [dev] - XXX. XX, XXXX ### Added +* Added `DPCTLQueue_GetLastSubmitError` C-API function, which returns the reason why the last kernel submission in the calling thread failed [gh-XXXX](https://github.com/IntelPython/dpctl/pull/XXXX) ### Changed @@ -16,6 +17,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 * Removed the deprecated `dpctl.program` submodule, use `dpctl.compiler` instead [gh-2407](https://github.com/IntelPython/dpctl/pull/2407) ### Fixed +* `dpctl.SyclKernelSubmitError` raised by `dpctl.SyclQueue.submit` and `dpctl.SyclQueue.submit_async` now includes the reason reported by the SYCL runtime, and the `dpctl.compiler.create_kernel_bundle_from_sycl_source` docstring documents the `-fno-sycl-id-queries-fit-in-int` option needed with DPC++ 2026.2 to launch kernels over more than `INT_MAX` work-items [gh-XXXX](https://github.com/IntelPython/dpctl/pull/XXXX) * Removed stale `select_host_device` and `has_host_device` names from `dpctl.__all__`, which made `from dpctl import *` fail with `AttributeError` [gh-2410](https://github.com/IntelPython/dpctl/pull/2410) ### Maintenance diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index 3708ea9130..a3abff4d70 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -649,6 +649,7 @@ cdef extern from "syclinterface/dpctl_sycl_queue_interface.h": size_t NDims, const DPCTLSyclEventRef *DepEvents, size_t NDepEvents) + cdef const char *DPCTLQueue_GetLastSubmitError() cdef void DPCTLQueue_Wait(const DPCTLSyclQueueRef QRef) nogil cdef DPCTLSyclEventRef DPCTLQueue_Memcpy( const DPCTLSyclQueueRef Q, diff --git a/dpctl/_sycl_queue.pyx b/dpctl/_sycl_queue.pyx index 3fbf9c7bf6..ad0b8bb6c9 100644 --- a/dpctl/_sycl_queue.pyx +++ b/dpctl/_sycl_queue.pyx @@ -53,6 +53,7 @@ from ._backend cimport ( # noqa: E211 DPCTLQueue_GetBackend, DPCTLQueue_GetContext, DPCTLQueue_GetDevice, + DPCTLQueue_GetLastSubmitError, DPCTLQueue_HasEnableProfiling, DPCTLQueue_Hash, DPCTLQueue_IsInOrder, @@ -1459,7 +1460,8 @@ cdef class SyclQueue(_SyclQueue): of arguments the kernel takes. SyclKernelSubmitError: If ``kernel`` does not reference a ``sycl::kernel``, or if - submission to the queue fails. + submission to the queue fails. The message includes the + reason reported by the SYCL runtime, if any. """ cdef void **kargs = NULL cdef _arg_data_type *kargty = NULL @@ -1591,8 +1593,13 @@ cdef class SyclQueue(_SyclQueue): free(depEvents) if Eref is NULL: + # the C API keeps the reason, e.g. the text of the SYCL exception + reason = DPCTLQueue_GetLastSubmitError().decode( + "utf-8", errors="replace" + ) + msg = "Kernel submission to Sycl queue failed" raise SyclKernelSubmitError( - "Kernel submission to Sycl queue failed." + f"{msg}: {reason}" if reason else f"{msg}." ) return SyclEvent._create(Eref) @@ -1637,7 +1644,8 @@ cdef class SyclQueue(_SyclQueue): of arguments the kernel takes. SyclKernelSubmitError: If ``kernel`` does not reference a ``sycl::kernel``, or if - submission to the queue fails. + submission to the queue fails. The message includes the + reason reported by the SYCL runtime, if any. """ cdef SyclEvent e = self.submit_async(kernel, args, gS, lS, dEvents) e.wait() diff --git a/dpctl/compiler/_compiler.pyx b/dpctl/compiler/_compiler.pyx index 58d7a56310..fa1bd93540 100644 --- a/dpctl/compiler/_compiler.pyx +++ b/dpctl/compiler/_compiler.pyx @@ -720,6 +720,16 @@ cpdef create_kernel_bundle_from_sycl_source(SyclQueue q, when compiling the program. Default: ``None`` + .. note:: + By default, kernels are compiled under the assumption + that every global ID, range and offset fits into an + ``int``. Starting with DPC++ 2026.2 the runtime checks + this at each launch, so submitting such a kernel over + more than ``INT_MAX`` (2,147,483,647) work-items raises + :class:`dpctl.SyclKernelSubmitError`. Pass + ``"-fno-sycl-id-queries-fit-in-int"`` in ``copts`` to + compile kernels that are launched over larger ranges. + Returns: kernel_bundle (:class:`.SyclKernelBundle`) A :class:`.SyclKernelBundle` object wrapping the diff --git a/dpctl/tests/test_sycl_kernel_submit.py b/dpctl/tests/test_sycl_kernel_submit.py index e173035ae6..d0db4ac7e5 100644 --- a/dpctl/tests/test_sycl_kernel_submit.py +++ b/dpctl/tests/test_sycl_kernel_submit.py @@ -368,3 +368,28 @@ def test_submit_wrong_number_of_args(): [x_usm, x_usm, x_usm, ctypes.c_int(2), ctypes.c_int(3)], [n], ) + + +def test_submit_error_message(): + try: + q = dpctl.SyclQueue() + except dpctl.SyclQueueCreationError: + pytest.skip("Could not create default queue") + fn = get_spirv_abspath("multi_kernel.spv") + with open(fn, "br") as f: + spirv_bytes = f.read() + kb = dpc.create_kernel_bundle_from_spirv(q, spirv_bytes) + krn = kb.get_sycl_kernel("axpy") + + n = 16 + x_usm = dpm.MemoryUSMDevice(n * 4, queue=q) + # the SYCL runtime rejects a global range that is not a multiple of + # the local range, and its reason must reach the Python exception + prefix = "Kernel submission to Sycl queue failed: " + with pytest.raises(dpctl.SyclKernelSubmitError) as excinfo: + q.submit_async( + krn, [x_usm, x_usm, x_usm, ctypes.c_int(2)], [n + 1], [n] + ) + msg = str(excinfo.value) + assert msg.startswith(prefix) + assert len(msg) > len(prefix) diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h index fb8b76a236..21c65f7aa1 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h @@ -276,6 +276,23 @@ DPCTLQueue_SubmitNDRange(__dpctl_keep const DPCTLSyclKernelRef KRef, __dpctl_keep const DPCTLSyclEventRef *DepEvents, size_t NDepEvents); +/*! + * @brief Returns the reason why the last call to DPCTLQueue_SubmitRange or + * DPCTLQueue_SubmitNDRange made in the calling thread failed. + * + * Both functions return NULL on failure and store the reason, e.g. the + * ``what()`` text of the SYCL exception thrown by ``sycl::queue::submit``. + * A successful call clears it. + * + * @return A NUL-terminated string, empty if the last call succeeded or no + * call has been made in this thread. The string is owned by + * libsyclinterface and is valid until the next call to either + * submit function in the same thread. + * @ingroup QueueInterface + */ +DPCTL_API +const char *DPCTLQueue_GetLastSubmitError(void); + /*! * @brief Calls the ``sycl::queue::submit`` function to do a blocking wait on * all enqueued tasks in the queue. diff --git a/libsyclinterface/source/dpctl_sycl_queue_interface.cpp b/libsyclinterface/source/dpctl_sycl_queue_interface.cpp index a7d2a49c8b..860efee1e1 100644 --- a/libsyclinterface/source/dpctl_sycl_queue_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_queue_interface.cpp @@ -39,6 +39,7 @@ #include #include #include +#include #include /* SYCL headers */ #include #include @@ -120,6 +121,22 @@ using namespace sycl; } \ } while (0); +namespace +{ +// Reason for the last failed DPCTLQueue_SubmitRange or +// DPCTLQueue_SubmitNDRange call in the calling thread, or an empty string if +// the last call succeeded. +thread_local std::string last_submit_error; + +std::nullptr_t +submit_failed(const std::string &Msg, const char *func_name, int line_num) +{ + last_submit_error = Msg; + error_handler(Msg, __FILE__, func_name, line_num, error_level::error); + return nullptr; +} +} // namespace + namespace { static_assert(__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_VERSION_REQUIRED, @@ -498,16 +515,15 @@ DPCTLQueue_SubmitRange(__dpctl_keep const DPCTLSyclKernelRef KRef, __dpctl_keep const DPCTLSyclEventRef *DepEvents, size_t NDepEvents) { - if (!KRef) { - error_handler("Cannot submit the kernel as the kernel is a nullptr.", - __FILE__, __func__, __LINE__, error_level::error); - return nullptr; - } - if (!QRef) { - error_handler("Cannot submit the kernel as the queue is a nullptr.", - __FILE__, __func__, __LINE__, error_level::error); - return nullptr; - } + last_submit_error.clear(); + if (!KRef) + return submit_failed( + "Cannot submit the kernel as the kernel is a nullptr.", __func__, + __LINE__); + if (!QRef) + return submit_failed( + "Cannot submit the kernel as the queue is a nullptr.", __func__, + __LINE__); auto Kernel = unwrap(KRef); auto Queue = unwrap(QRef); @@ -547,18 +563,15 @@ DPCTLQueue_SubmitRange(__dpctl_keep const DPCTLSyclKernelRef KRef, return wrap(new event(std::move(e))); } default: - error_handler("Range cannot be greater than three " - "dimensions.", - __FILE__, __func__, __LINE__, error_level::error); - return nullptr; + return submit_failed("Range cannot be greater than three " + "dimensions.", + __func__, __LINE__); } } catch (std::exception const &e) { - error_handler(e, __FILE__, __func__, __LINE__, error_level::error); - return nullptr; + return submit_failed(e.what(), __func__, __LINE__); } catch (...) { - error_handler("Unknown exception encountered", __FILE__, __func__, - __LINE__, error_level::error); - return nullptr; + return submit_failed("Unknown exception encountered", __func__, + __LINE__); } } @@ -574,16 +587,15 @@ DPCTLQueue_SubmitNDRange(__dpctl_keep const DPCTLSyclKernelRef KRef, __dpctl_keep const DPCTLSyclEventRef *DepEvents, size_t NDepEvents) { - if (!KRef) { - error_handler("Cannot submit the kernel as the kernel is a nullptr.", - __FILE__, __func__, __LINE__, error_level::error); - return nullptr; - } - if (!QRef) { - error_handler("Cannot submit the kernel as the queue is a nullptr.", - __FILE__, __func__, __LINE__, error_level::error); - return nullptr; - } + last_submit_error.clear(); + if (!KRef) + return submit_failed( + "Cannot submit the kernel as the kernel is a nullptr.", __func__, + __LINE__); + if (!QRef) + return submit_failed( + "Cannot submit the kernel as the queue is a nullptr.", __func__, + __LINE__); auto Kernel = unwrap(KRef); auto Queue = unwrap(QRef); @@ -627,21 +639,23 @@ DPCTLQueue_SubmitNDRange(__dpctl_keep const DPCTLSyclKernelRef KRef, return wrap(new event(std::move(e))); } default: - error_handler("Range cannot be greater than three " - "dimensions.", - __FILE__, __func__, __LINE__, error_level::error); - return nullptr; + return submit_failed("Range cannot be greater than three " + "dimensions.", + __func__, __LINE__); } } catch (std::exception const &e) { - error_handler(e, __FILE__, __func__, __LINE__, error_level::error); - return nullptr; + return submit_failed(e.what(), __func__, __LINE__); } catch (...) { - error_handler("Unknown exception encountered", __FILE__, __func__, - __LINE__, error_level::error); - return nullptr; + return submit_failed("Unknown exception encountered", __func__, + __LINE__); } } +const char *DPCTLQueue_GetLastSubmitError(void) +{ + return last_submit_error.c_str(); +} + void DPCTLQueue_Wait(__dpctl_keep DPCTLSyclQueueRef QRef) { // \todo what happens if the QRef is null or a pointer to a valid sycl diff --git a/libsyclinterface/tests/test_sycl_queue_submit.cpp b/libsyclinterface/tests/test_sycl_queue_submit.cpp index fc0611a548..69382313fa 100644 --- a/libsyclinterface/tests/test_sycl_queue_submit.cpp +++ b/libsyclinterface/tests/test_sycl_queue_submit.cpp @@ -400,11 +400,56 @@ TEST_F(TestQueueSubmit, CheckForNullKernel) ERef = DPCTLQueue_SubmitRange(nullptr, QRef, nullptr, nullptr, 0, Range, RANGE_NDIMS, nullptr, 0)); ASSERT_TRUE(ERef == nullptr); + EXPECT_STRNE(DPCTLQueue_GetLastSubmitError(), ""); EXPECT_NO_FATAL_FAILURE(ERef = DPCTLQueue_SubmitNDRange( nullptr, QRef, nullptr, nullptr, 0, Range, lRange, RANGE_NDIMS, nullptr, 0)); ASSERT_TRUE(ERef == nullptr); + EXPECT_STRNE(DPCTLQueue_GetLastSubmitError(), ""); +} + +TEST_F(TestQueueSubmit, CheckLastSubmitError) +{ + // a global range that is not a multiple of the local range makes + // sycl::queue::submit throw + std::size_t gRange[] = {SIZE + 1}; + std::size_t lRange[] = {SIZE}; + std::size_t RANGE_NDIMS = 1; + DPCTLSyclEventRef ERef = nullptr; + + auto kernel = DPCTLKernelBundle_GetKernel(KBRef, "_ZTS11RangeKernelIiE"); + ASSERT_TRUE(kernel != nullptr); + + auto a = DPCTLmalloc_shared(SIZE * sizeof(int), QRef); + auto b = DPCTLmalloc_shared(SIZE * sizeof(int), QRef); + auto c = DPCTLmalloc_shared(SIZE * sizeof(int), QRef); + ASSERT_TRUE(a && b && c); + int scalarVal = 3; + void *args[] = {unwrap(a), unwrap(b), unwrap(c), + (void *)&scalarVal}; + DPCTLKernelArgType argTypes[] = {DPCTL_VOID_PTR, DPCTL_VOID_PTR, + DPCTL_VOID_PTR, DPCTL_INT32_T}; + + EXPECT_NO_FATAL_FAILURE( + ERef = DPCTLQueue_SubmitNDRange(kernel, QRef, args, argTypes, 4, gRange, + lRange, RANGE_NDIMS, nullptr, 0)); + ASSERT_TRUE(ERef == nullptr); + EXPECT_STRNE(DPCTLQueue_GetLastSubmitError(), ""); + + // a successful submission clears the error + EXPECT_NO_FATAL_FAILURE( + ERef = DPCTLQueue_SubmitRange(kernel, QRef, args, argTypes, 4, lRange, + RANGE_NDIMS, nullptr, 0)); + ASSERT_TRUE(ERef != nullptr); + EXPECT_STREQ(DPCTLQueue_GetLastSubmitError(), ""); + EXPECT_NO_FATAL_FAILURE(DPCTLEvent_Wait(ERef)); + EXPECT_NO_FATAL_FAILURE(DPCTLEvent_Delete(ERef)); + + EXPECT_NO_FATAL_FAILURE(DPCTLfree_with_queue(a, QRef)); + EXPECT_NO_FATAL_FAILURE(DPCTLfree_with_queue(b, QRef)); + EXPECT_NO_FATAL_FAILURE(DPCTLfree_with_queue(c, QRef)); + EXPECT_NO_FATAL_FAILURE(DPCTLKernel_Delete(kernel)); } TEST_F(TestQueueSubmit, CheckForNullQueue) From a718e1898aedf18d082aba712a06748c89a16b81 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Wed, 4 Nov 2026 13:11:37 +0100 Subject: [PATCH 2/3] Add gh-2414 to the changelog --- CHANGELOG.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 3b3c31e442..70d2b16bd1 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -7,7 +7,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 ## [dev] - XXX. XX, XXXX ### Added -* Added `DPCTLQueue_GetLastSubmitError` C-API function, which returns the reason why the last kernel submission in the calling thread failed [gh-XXXX](https://github.com/IntelPython/dpctl/pull/XXXX) +* Added `DPCTLQueue_GetLastSubmitError` C-API function, which returns the reason why the last kernel submission in the calling thread failed [gh-2414](https://github.com/IntelPython/dpctl/pull/2414) ### Changed @@ -17,7 +17,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 * Removed the deprecated `dpctl.program` submodule, use `dpctl.compiler` instead [gh-2407](https://github.com/IntelPython/dpctl/pull/2407) ### Fixed -* `dpctl.SyclKernelSubmitError` raised by `dpctl.SyclQueue.submit` and `dpctl.SyclQueue.submit_async` now includes the reason reported by the SYCL runtime, and the `dpctl.compiler.create_kernel_bundle_from_sycl_source` docstring documents the `-fno-sycl-id-queries-fit-in-int` option needed with DPC++ 2026.2 to launch kernels over more than `INT_MAX` work-items [gh-XXXX](https://github.com/IntelPython/dpctl/pull/XXXX) +* `dpctl.SyclKernelSubmitError` raised by `dpctl.SyclQueue.submit` and `dpctl.SyclQueue.submit_async` now includes the reason reported by the SYCL runtime, and the `dpctl.compiler.create_kernel_bundle_from_sycl_source` docstring documents the `-fno-sycl-id-queries-fit-in-int` option needed with DPC++ 2026.2 to launch kernels over more than `INT_MAX` work-items [gh-2414](https://github.com/IntelPython/dpctl/pull/2414) * Removed stale `select_host_device` and `has_host_device` names from `dpctl.__all__`, which made `from dpctl import *` fail with `AttributeError` [gh-2410](https://github.com/IntelPython/dpctl/pull/2410) ### Maintenance From dad43718507bfec4b4d4395c91d6b0ebf8856127 Mon Sep 17 00:00:00 2001 From: Anton Volkov Date: Wed, 4 Nov 2026 13:42:42 +0100 Subject: [PATCH 3/3] Use an oversized work-group to make kernel submission fail in tests A global range that is not a multiple of the local range is accepted by OpenCL devices supporting non-uniform work-groups, e.g. the OpenCL CPU device used in CI, so the submission did not fail there. A local range twice the device's maximum work-group size is rejected on every device. --- dpctl/tests/test_sycl_kernel_submit.py | 11 ++++++----- .../tests/test_sycl_queue_submit.cpp | 19 ++++++++++++++----- 2 files changed, 20 insertions(+), 10 deletions(-) diff --git a/dpctl/tests/test_sycl_kernel_submit.py b/dpctl/tests/test_sycl_kernel_submit.py index d0db4ac7e5..34e8bebef9 100644 --- a/dpctl/tests/test_sycl_kernel_submit.py +++ b/dpctl/tests/test_sycl_kernel_submit.py @@ -383,13 +383,14 @@ def test_submit_error_message(): n = 16 x_usm = dpm.MemoryUSMDevice(n * 4, queue=q) - # the SYCL runtime rejects a global range that is not a multiple of - # the local range, and its reason must reach the Python exception + # the SYCL runtime rejects a local range that exceeds the device's + # maximum work-group size, and its reason must reach the Python + # exception; a global range that is not a multiple of the local range + # is not rejected, as OpenCL devices may support non-uniform work-groups + wg = 2 * q.sycl_device.max_work_group_size prefix = "Kernel submission to Sycl queue failed: " with pytest.raises(dpctl.SyclKernelSubmitError) as excinfo: - q.submit_async( - krn, [x_usm, x_usm, x_usm, ctypes.c_int(2)], [n + 1], [n] - ) + q.submit_async(krn, [x_usm, x_usm, x_usm, ctypes.c_int(2)], [wg], [wg]) msg = str(excinfo.value) assert msg.startswith(prefix) assert len(msg) > len(prefix) diff --git a/libsyclinterface/tests/test_sycl_queue_submit.cpp b/libsyclinterface/tests/test_sycl_queue_submit.cpp index 69382313fa..7b3b2a4831 100644 --- a/libsyclinterface/tests/test_sycl_queue_submit.cpp +++ b/libsyclinterface/tests/test_sycl_queue_submit.cpp @@ -411,10 +411,19 @@ TEST_F(TestQueueSubmit, CheckForNullKernel) TEST_F(TestQueueSubmit, CheckLastSubmitError) { - // a global range that is not a multiple of the local range makes - // sycl::queue::submit throw - std::size_t gRange[] = {SIZE + 1}; - std::size_t lRange[] = {SIZE}; + // a local range that exceeds the device's maximum work-group size makes + // sycl::queue::submit throw; a global range that is not a multiple of + // the local range does not, as OpenCL devices may support non-uniform + // work-groups + auto DRef = DPCTLQueue_GetDevice(QRef); + ASSERT_TRUE(DRef != nullptr); + std::size_t maxWGSize = DPCTLDevice_GetMaxWorkGroupSize(DRef); + EXPECT_NO_FATAL_FAILURE(DPCTLDevice_Delete(DRef)); + ASSERT_TRUE(maxWGSize > 0); + + std::size_t gRange[] = {2 * maxWGSize}; + std::size_t lRange[] = {2 * maxWGSize}; + std::size_t Range[] = {SIZE}; std::size_t RANGE_NDIMS = 1; DPCTLSyclEventRef ERef = nullptr; @@ -439,7 +448,7 @@ TEST_F(TestQueueSubmit, CheckLastSubmitError) // a successful submission clears the error EXPECT_NO_FATAL_FAILURE( - ERef = DPCTLQueue_SubmitRange(kernel, QRef, args, argTypes, 4, lRange, + ERef = DPCTLQueue_SubmitRange(kernel, QRef, args, argTypes, 4, Range, RANGE_NDIMS, nullptr, 0)); ASSERT_TRUE(ERef != nullptr); EXPECT_STREQ(DPCTLQueue_GetLastSubmitError(), "");