diff --git a/CHANGELOG.md b/CHANGELOG.md index 4b3d0c90d1..70d2b16bd1 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-2414](https://github.com/IntelPython/dpctl/pull/2414) ### 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-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 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..34e8bebef9 100644 --- a/dpctl/tests/test_sycl_kernel_submit.py +++ b/dpctl/tests/test_sycl_kernel_submit.py @@ -368,3 +368,29 @@ 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 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)], [wg], [wg]) + 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..7b3b2a4831 100644 --- a/libsyclinterface/tests/test_sycl_queue_submit.cpp +++ b/libsyclinterface/tests/test_sycl_queue_submit.cpp @@ -400,11 +400,65 @@ 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 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; + + 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, Range, + 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)