Pass the SYCL submission error to SyclKernelSubmitError by antonwolfy · Pull Request #2414 · IntelPython/dpctl

antonwolfy · 2026-11-04T12:11:31Z

Kernels built by dpctl.compiler.create_kernel_bundle_from_sycl_source get the default int ID range. With DPC++ 2026.2 the runtime checks it at each launch, so a launch over more than INT_MAX work-items fails unless -fno-sycl-id-queries-fit-in-int is passed in copts. From Python the failure looked like this:

SyclKernelSubmitError: Kernel submission to Sycl queue failed.

The reason was printed only with DPCTL_VERBOSITY=warning, because DPCTLQueue_SubmitRange and DPCTLQueue_SubmitNDRange catch the SYCL exception and return NULL. Other submission errors were hidden the same way.

This PR:

keeps the reason for a failed DPCTLQueue_SubmitRange/DPCTLQueue_SubmitNDRange call in a thread-local string and adds the DPCTLQueue_GetLastSubmitError C-API function to read it. A successful call clears it. The existing function signatures don't change.

appends that reason to the SyclKernelSubmitError message raised by SyclQueue.submit_async and SyclQueue.submit. The same launch now raises:

SyclKernelSubmitError: Kernel submission to Sycl queue failed: The kernel was compiled with -fsycl-id-queries-range=int, but the provided range/offset exceeds the maximum value storable in an int. Either reduce the range/offset or recompile the kernel with -fsycl-id-queries-range=[uint|size_t].

documents -fno-sycl-id-queries-fit-in-int under copts in the create_kernel_bundle_from_sycl_source docstring.

Tested with DPC++ 2026.2 on OpenCL CPU and Level Zero GPU: an nd_range launch over 2,147,483,712 work-items raises the message above without the option, and runs with the right result with it. New tests: the C-API test TestQueueSubmit.CheckLastSubmitError and the Python test test_submit_error_message, which submit an nd_range whose global range is not a multiple of the local range. The full test suites pass: C-API tests 2505 passed, Python 3178 passed, 1120 skipped, 1 xfailed.

Have you provided a meaningful PR description?
Have you added a test, reproducer or referred to an issue with a reproducer?
Have you tested your changes locally for CPU and GPU devices?
Have you made sure that new changes do not introduce compiler warnings?
Have you checked performance impact of proposed changes?
Have you added documentation for your changes, if necessary?
Have you added your changes to the changelog?
If this PR is a work in progress, are you opening the PR as a draft?

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Copilot

🟢 Approval recommended

No active review issues are recorded for the change that exposes SYCL submission failure details in Python exceptions.

No open issues 🎉

What changed in this PR

This PR carries SYCL kernel submission failure details from the C API into Python exceptions, making launch errors diagnosable without enabling verbose logging.

Changes:

Adds thread-local submission error storage and a C API accessor, clearing the error on successful submission.
Includes the runtime error reason in SyclKernelSubmitError and adds C API and Python regression tests.
Documents the compiler option for larger launch ranges and records the changes in the changelog.

File	Description
`libsyclinterface/tests/test_sycl_queue_submit.cpp`	Tests error retrieval and clearing after successful submission.
`libsyclinterface/source/dpctl_sycl_queue_interface.cpp`	Stores submission errors per thread and implements the accessor.
`libsyclinterface/include/syclinterface/dpctl_sycl_queue_interface.h`	Declares the accessor and documents string ownership and lifetime.
`dpctl/tests/test_sycl_kernel_submit.py`	Checks that submission exceptions include an error reason.
`dpctl/compiler/_compiler.pyx`	Documents the compiler option for larger launch ranges.
`dpctl/_sycl_queue.pyx`	Appends submission error details to Python exceptions.
`dpctl/_backend.pxd`	Exposes the new C API accessor to Cython.
`CHANGELOG.md`	Records the new accessor, improved exceptions, and documentation.

_{Review effort: Balanced}

github-actions · 2026-11-04T12:15:05Z

View rendered docs @ https://intelpython.github.io/dpctl/pulls/2414/index.html

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.

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.

coveralls · 2026-11-04T12:48:28Z

coverage: 74.967% (+0.07%) from 74.893% — kernel-submit-error-message into master

ndgrigorian · 2026-11-05T05:07:26Z

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


These are two separate changes, can we split them into two entries?

Suggested change

* `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)

* `dpctl.SyclKernelSubmitError` raised by `dpctl.SyclQueue.submit` and `dpctl.SyclQueue.submit_async` now includes the reason reported by the SYCL runtime [gh-2414](https://github.com/IntelPython/dpctl/pull/2414)

* `dpctl.compiler.create_kernel_bundle_from_sycl_source` docstring now 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)

ndgrigorian · 2026-11-05T05:07:26Z

 #include <exception>
 #include <sstream>
 #include <stdexcept>
+#include <string>


std::nullptr_t is from <cstddef>, which only comes in through <string> transitively. Might as well include it

Suggested change

#include <string>

#include <cstddef>

#include <string>

ndgrigorian · 2026-11-05T05:07:26Z

+std::nullptr_t
+submit_failed(const std::string &Msg, const char *func_name, int line_num)
+{
+    last_submit_error = Msg;


This runs inside the catch handlers, so a bad_alloc here would escape through the extern "C" functions. Not new, error_handler already builds a stringstream in the same spot, but a try/catch (...) around the assignment would cover it

ndgrigorian · 2026-11-05T05:07:26Z

+    EXPECT_NO_FATAL_FAILURE(
+        ERef = DPCTLQueue_SubmitNDRange(kernel, QRef, args, argTypes, 4, gRange,
+                                        lRange, RANGE_NDIMS, nullptr, 0));
+    ASSERT_TRUE(ERef == nullptr);


If this or the ASSERT_TRUE(ERef != nullptr) below fails, the test returns early and leaks a, b, c and the kernel. Only matters when the test is already failing, and other tests here do the same, so minor

ndgrigorian · 2026-11-05T05:07:27Z


        if Eref is NULL:
+            # the C API keeps the reason, e.g. the text of the SYCL exception
+            reason = DPCTLQueue_GetLastSubmitError().decode(


Nothing between the submit and this read touches the thread-local now, but reading the error right after DPCTLQueue_Submit* (before the free calls) would make that ordering explicit

antonwolfy requested review from ndgrigorian and vlad-perevezentsev as code owners November 4, 2026 12:11

antonwolfy added a commit that referenced this pull request Nov 4, 2026

Add gh-2414 to the changelog

d9f2ded

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

antonwolfy self-assigned this Nov 4, 2026

antonwolfy requested a balanced review from Copilot November 4, 2026 12:12

Copilot started reviewing on behalf of antonwolfy November 4, 2026 12:13

View session

Copilot AI reviewed Nov 4, 2026

View reviewed changes

antonwolfy added 2 commits November 4, 2026 13:16

Add gh-2414 to the changelog

a718e18

antonwolfy force-pushed the kernel-submit-error-message branch from d9f2ded to a718e18 Compare November 4, 2026 12:16

ndgrigorian reviewed Nov 5, 2026

View reviewed changes

Navigation Menu

Repository navigation

Pass the SYCL submission error to SyclKernelSubmitError - #2414

Pass the SYCL submission error to SyclKernelSubmitError#2414
antonwolfy wants to merge 3 commits into
masterfrom
kernel-submit-error-message

antonwolfy commented Nov 4, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

github-actions Bot commented Nov 4, 2026

Uh oh!

coveralls commented Nov 4, 2026

Uh oh!

ndgrigorian Nov 5, 2026

Uh oh!

ndgrigorian Nov 5, 2026

Uh oh!

ndgrigorian Nov 5, 2026

Uh oh!

ndgrigorian Nov 5, 2026

Uh oh!

ndgrigorian Nov 5, 2026

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

	* `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)
	* `dpctl.SyclKernelSubmitError` raised by `dpctl.SyclQueue.submit` and `dpctl.SyclQueue.submit_async` now includes the reason reported by the SYCL runtime [gh-2414](https://github.com/IntelPython/dpctl/pull/2414)
	* `dpctl.compiler.create_kernel_bundle_from_sycl_source` docstring now 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)

Conversation

antonwolfy commented Nov 4, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟢 Approval recommended

Uh oh!

github-actions Bot commented Nov 4, 2026

Uh oh!

coveralls commented Nov 4, 2026

Uh oh!

ndgrigorian Nov 5, 2026

Choose a reason for hiding this comment

Uh oh!

ndgrigorian Nov 5, 2026

Choose a reason for hiding this comment

Uh oh!

ndgrigorian Nov 5, 2026

Choose a reason for hiding this comment

Uh oh!

ndgrigorian Nov 5, 2026

Choose a reason for hiding this comment

Uh oh!

ndgrigorian Nov 5, 2026

Choose a reason for hiding this comment

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants