[SYCL][HIP][CUDA] Report errc::nd_range for oversized work-group size

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
   }

sycl/test-e2e/Basic/gpu_max_wgs_error.cpp

@@ -1,8 +1,18 @@
 // REQUIRES: gpu
-// UNSUPPORTED: cuda
 
-// UNSUPPORTED: hip
+// 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
 

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 <sycl/detail/core.hpp>
 
 int main() {