From 50c7a75594464542b237e2591caeec65a40f415f Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Sun, 19 Jul 2026 10:51:48 -0700 Subject: [PATCH 1/3] [SYCL][HIP][CUDA] Report errc::nd_range for oversized work-group size handleInvalidWorkGroupSize only produced the expected "Total number of work-items in a work-group cannot exceed ..." diagnostic (errc::nd_range) for OpenCL. On HIP and CUDA the launch was rejected by the driver but the runtime surfaced a generic error, so Basic/gpu_max_wgs_error.cpp failed. Handle the HIP and CUDA backends explicitly and throw the same errc::nd_range exception when the requested work-group size exceeds the device maximum. Re-enables Basic/gpu_max_wgs_error.cpp on HIP. Co-authored-by: Cursor --- sycl/source/detail/error_handling/error_handling.cpp | 11 +++++++++++ sycl/test-e2e/Basic/gpu_max_wgs_error.cpp | 2 -- 2 files changed, 11 insertions(+), 2 deletions(-) diff --git a/sycl/source/detail/error_handling/error_handling.cpp b/sycl/source/detail/error_handling/error_handling.cpp index 8d263526e8209..14b5d5a76ecbb 100644 --- a/sycl/source/detail/error_handling/error_handling.cpp +++ b/sycl/source/detail/error_handling/error_handling.cpp @@ -240,6 +240,17 @@ void handleInvalidWorkGroupSize(const device_impl &DeviceImpl, make_error_code(errc::nd_range), "Total number of work-items in a work-group cannot exceed " + std::to_string(KernelWGSize) + " for this kernel"); + } else if (Backend == sycl::backend::ext_oneapi_cuda || + Backend == sycl::backend::ext_oneapi_hip) { + // CUDA and HIP: + // The underlying driver rejects launches whose total work-group size + // exceeds UR_DEVICE_INFO_MAX_WORK_GROUP_SIZE, so surface the same + // diagnostic as the OpenCL 1.x path above. + if (TotalNumberOfWIs > MaxWGSize) + throw sycl::exception( + make_error_code(errc::nd_range), + "Total number of work-items in a work-group cannot exceed " + + std::to_string(MaxWGSize)); } else { // TODO: Should probably have something similar for the other backends } diff --git a/sycl/test-e2e/Basic/gpu_max_wgs_error.cpp b/sycl/test-e2e/Basic/gpu_max_wgs_error.cpp index 968167f6d8f3f..88340ad8275b9 100644 --- a/sycl/test-e2e/Basic/gpu_max_wgs_error.cpp +++ b/sycl/test-e2e/Basic/gpu_max_wgs_error.cpp @@ -1,8 +1,6 @@ // REQUIRES: gpu // UNSUPPORTED: cuda -// UNSUPPORTED: hip -// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/22300 // RUN: %{build} -o %t.out -fno-sycl-id-queries-fit-in-int // RUN: %{run} %t.out From 9eeb34d2f7fef07079c0bfdea30bb8423f14eeff Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Sun, 19 Jul 2026 09:22:00 -0700 Subject: [PATCH 2/3] [SYCL][E2E] Re-enable Regression/invalid_reqd_wg_size_correct_exception on HIP This test passes on HIP (gfx942) with the current toolchain, so remove the UNSUPPORTED: hip marker tracked by intel/llvm#22300. Co-authored-by: Cursor --- .../Regression/invalid_reqd_wg_size_correct_exception.cpp | 3 --- 1 file changed, 3 deletions(-) diff --git a/sycl/test-e2e/Regression/invalid_reqd_wg_size_correct_exception.cpp b/sycl/test-e2e/Regression/invalid_reqd_wg_size_correct_exception.cpp index a916f65f94e52..4c949a59134b8 100644 --- a/sycl/test-e2e/Regression/invalid_reqd_wg_size_correct_exception.cpp +++ b/sycl/test-e2e/Regression/invalid_reqd_wg_size_correct_exception.cpp @@ -1,9 +1,6 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -// UNSUPPORTED: hip -// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/22300 - #include int main() { From 97a1862c8012dc1eec88807aac3afbad30d2a632 Mon Sep 17 00:00:00 2001 From: Zheming Jin Date: Mon, 20 Jul 2026 09:02:41 -0700 Subject: [PATCH 3/3] [SYCL][E2E] Document why gpu_max_wgs_error stays UNSUPPORTED on CUDA The CUDA backend does report errc::nd_range for an oversized work-group, but for the work-group size this test uses (max_work_item_sizes in every dimension) the CUDA driver rejects the launch as UR_RESULT_ERROR_OUT_OF_RESOURCES (register exhaustion), which is handled by handleOutOfResources() and throws an nd_range exception with a different, register-count message. Verified on an H100 that the "Total number of work-items in a work-group cannot exceed" branch added in handleInvalidWorkGroupSize() is reached (and matches) only when each dimension is within limits but their product exceeds the max work-group size. Add an explanation and tracker so the CUDA exclusion is documented. Co-authored-by: Cursor --- sycl/test-e2e/Basic/gpu_max_wgs_error.cpp | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/sycl/test-e2e/Basic/gpu_max_wgs_error.cpp b/sycl/test-e2e/Basic/gpu_max_wgs_error.cpp index 88340ad8275b9..4097ab6306c0e 100644 --- a/sycl/test-e2e/Basic/gpu_max_wgs_error.cpp +++ b/sycl/test-e2e/Basic/gpu_max_wgs_error.cpp @@ -1,5 +1,17 @@ // REQUIRES: gpu + +// The runtime does throw errc::nd_range for an oversized work-group on CUDA, +// but for the work-group size used here (max_work_item_sizes in every +// dimension) the CUDA driver rejects the launch as +// UR_RESULT_ERROR_OUT_OF_RESOURCES (register exhaustion) rather than +// UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE. That is handled by +// handleOutOfResources(), which throws an nd_range exception with a different, +// register-count message. The "Total number of work-items in a work-group +// cannot exceed" branch added in handleInvalidWorkGroupSize() is only reached +// when each dimension is within limits but their product exceeds the maximum +// work-group size, so this generic test cannot assert that message on CUDA. // UNSUPPORTED: cuda +// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/22300 // RUN: %{build} -o %t.out -fno-sycl-id-queries-fit-in-int // RUN: %{run} %t.out