Pass the SYCL submission error to SyclKernelSubmitError

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

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,

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()

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

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)

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.

libsyclinterface/source/dpctl_sycl_queue_interface.cpp

@@ -39,6 +39,7 @@
 #include <exception>
 #include <sstream>
 #include <stdexcept>
+#include <string>
 #include <sycl/sycl.hpp> /* SYCL headers   */
 #include <utility>
 #include <vector>
@@ -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<kernel>(KRef);
     auto Queue = unwrap<queue>(QRef);
@@ -547,18 +563,15 @@ DPCTLQueue_SubmitRange(__dpctl_keep const DPCTLSyclKernelRef KRef,
             return wrap<event>(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<kernel>(KRef);
     auto Queue = unwrap<queue>(QRef);
@@ -627,21 +639,23 @@ DPCTLQueue_SubmitNDRange(__dpctl_keep const DPCTLSyclKernelRef KRef,
             return wrap<event>(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

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<void>(a), unwrap<void>(b), unwrap<void>(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)