[SYCL][CUDA] Improve kernel launch error handling for out-of-registers

sycl/plugins/unified_runtime/CMakeLists.txt

@@ -100,13 +100,13 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT)
   endfunction()
 
   set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit 9f783837089c970a22cda08f768aa3dbed38f0d3
-  # Merge: c015f892 b9442104
+  # commit 5083f4f96557672b7b6a55ea53347896d40549d7
+  # Merge: a97eed15 4c3f9abe
   # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
-  # Date:   Fri May 31 10:20:23 2024 +0100
-  #     Merge pull request #1533 from AllanZyne/sanitizer-buffer
-  #     [DeviceSanitizer] Support detecting out-of-bounds errors on sycl::buffer
-  set(UNIFIED_RUNTIME_TAG 9f783837089c970a22cda08f768aa3dbed38f0d3)
+  # Date:   Fri May 31 17:20:01 2024 +0100
+  #     Merge pull request #1397 from GeorgeWeb/georgi/check-allocation-error-on-event-from-native-handle
+  #     [CUDA][HIP] Catch and report bad_alloc errors for event object creation
+  set(UNIFIED_RUNTIME_TAG 5083f4f96557672b7b6a55ea53347896d40549d7)
 
   fetch_adapter_source(level_zero
     ${UNIFIED_RUNTIME_REPO}

sycl/source/detail/error_handling/error_handling.cpp

@@ -20,6 +20,60 @@ namespace sycl {
 inline namespace _V1 {
 namespace detail::enqueue_kernel_launch {
 
+void handleOutOfResources(const device_impl &DeviceImpl, pi_kernel Kernel,
+                          const NDRDescT &NDRDesc) {
+  sycl::platform Platform = DeviceImpl.get_platform();
+  sycl::backend Backend = Platform.get_backend();
+  if (Backend == sycl::backend::ext_oneapi_cuda) {
+    // PI_ERROR_OUT_OF_RESOURCES is returned when the kernel registers
+    // required for the launch config exceeds the maximum number of registers
+    // per block (PI_EXT_CODEPLAY_DEVICE_INFO_MAX_REGISTERS_PER_WORK_GROUP).
+    // This is if local_work_size[0] * ... * local_work_size[work_dim - 1]
+    // multiplied by PI_KERNEL_GROUP_INFO_NUM_REGS is greater than the value
+    // of PI_KERNEL_MAX_NUM_REGISTERS_PER_BLOCK. See Table 15: Technical
+    // Specifications per Compute Capability, for limitations.
+    const size_t TotalNumberOfWIs =
+        NDRDesc.LocalSize[0] * NDRDesc.LocalSize[1] * NDRDesc.LocalSize[2];
+
+    const uint32_t MaxRegistersPerBlock =
+        DeviceImpl.get_info<ext::codeplay::experimental::info::device::
+                                max_registers_per_work_group>();
+
+    const PluginPtr &Plugin = DeviceImpl.getPlugin();
+    sycl::detail::pi::PiDevice Device = DeviceImpl.getHandleRef();
+
+    uint32_t NumRegisters = 0;
+    Plugin->call<PiApiKind::piKernelGetGroupInfo>(
+        Kernel, Device, PI_KERNEL_GROUP_INFO_NUM_REGS, sizeof(NumRegisters),
+        &NumRegisters, nullptr);
+
+    const bool HasExceededAvailableRegisters =
+        TotalNumberOfWIs * NumRegisters > MaxRegistersPerBlock;
+
+    if (HasExceededAvailableRegisters) {
+      std::string message(
+          "Exceeded the number of registers available on the hardware.\n");
+      throw sycl::exception(
+          sycl::make_error_code(sycl::errc::nd_range),
+          // Additional information which can be helpful to the user.
+          message.append(
+              "\tThe number registers per work-group cannot exceed " +
+              std::to_string(MaxRegistersPerBlock) +
+              " for this kernel on this device.\n"
+              "\tThe kernel uses " +
+              std::to_string(NumRegisters) +
+              " registers per work-item for a total of " +
+              std::to_string(TotalNumberOfWIs) +
+              " work-items per work-group.\n"));
+    }
+  }
+  // Fallback
+  constexpr pi_result Error = PI_ERROR_OUT_OF_RESOURCES;
+  throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+                        "PI backend failed. PI backend returns:" +
+                            codeToString(Error));
+}
+
 void handleInvalidWorkGroupSize(const device_impl &DeviceImpl, pi_kernel Kernel,
                                 const NDRDescT &NDRDesc) {
   sycl::platform Platform = DeviceImpl.get_platform();
@@ -30,7 +84,6 @@ void handleInvalidWorkGroupSize(const device_impl &DeviceImpl, pi_kernel Kernel,
   bool IsOpenCLV1x = false;   // Backend is OpenCL 1.x
   bool IsOpenCLVGE20 = false; // Backend is Greater or Equal to OpenCL 2.0
   bool IsLevelZero = false;   // Backend is any OneAPI Level 0 version
-  bool IsCuda = false;        // Backend is CUDA
   auto Backend = Platform.get_backend();
   if (Backend == sycl::backend::opencl) {
     std::string VersionString =
@@ -41,8 +94,6 @@ void handleInvalidWorkGroupSize(const device_impl &DeviceImpl, pi_kernel Kernel,
         (VersionString.find("2.") == 0) || (VersionString.find("3.") == 0);
   } else if (Backend == sycl::backend::ext_oneapi_level_zero) {
     IsLevelZero = true;
-  } else if (Backend == sycl::backend::ext_oneapi_cuda) {
-    IsCuda = true;
   }
 
   const PluginPtr &Plugin = DeviceImpl.getPlugin();
@@ -243,46 +294,6 @@ void handleInvalidWorkGroupSize(const device_impl &DeviceImpl, pi_kernel Kernel,
           // else unknown.  fallback (below)
         }
       }
-    } else if (IsCuda) {
-      // CUDA:
-      // PI_ERROR_INVALID_WORK_GROUP_SIZE is returned when the kernel registers
-      // required for the launch config exceeds the maximum number of registers
-      // per block (PI_EXT_CODEPLAY_DEVICE_INFO_MAX_REGISTERS_PER_WORK_GROUP).
-      // This is if local_work_size[0] * ... * local_work_size[work_dim - 1]
-      // multiplied by PI_KERNEL_GROUP_INFO_NUM_REGS is greater than the value
-      // of PI_KERNEL_MAX_NUM_REGISTERS_PER_BLOCK. See Table 15: Technical
-      // Specifications per Compute Capability, for limitations.
-      const size_t TotalNumberOfWIs =
-          NDRDesc.LocalSize[0] * NDRDesc.LocalSize[1] * NDRDesc.LocalSize[2];
-
-      uint32_t NumRegisters = 0;
-      Plugin->call<PiApiKind::piKernelGetGroupInfo>(
-          Kernel, Device, PI_KERNEL_GROUP_INFO_NUM_REGS, sizeof(NumRegisters),
-          &NumRegisters, nullptr);
-
-      uint32_t MaxRegistersPerBlock =
-          DeviceImpl.get_info<ext::codeplay::experimental::info::device::
-                                  max_registers_per_work_group>();
-
-      const bool HasExceededAvailableRegisters =
-          TotalNumberOfWIs * NumRegisters > MaxRegistersPerBlock;
-
-      if (HasExceededAvailableRegisters) {
-        std::string message(
-            "Exceeded the number of registers available on the hardware.\n");
-        throw sycl::nd_range_error(
-            // Additional information which can be helpful to the user.
-            message.append(
-                "\tThe number registers per work-group cannot exceed " +
-                std::to_string(MaxRegistersPerBlock) +
-                " for this kernel on this device.\n"
-                "\tThe kernel uses " +
-                std::to_string(NumRegisters) +
-                " registers per work-item for a total of " +
-                std::to_string(TotalNumberOfWIs) +
-                " work-items per work-group.\n"),
-            PI_ERROR_INVALID_WORK_GROUP_SIZE);
-      }
     } else {
       // TODO: Decide what checks (if any) we need for the other backends
     }
@@ -352,6 +363,9 @@ void handleErrorOrWarning(pi_result Error, const device_impl &DeviceImpl,
   assert(Error != PI_SUCCESS &&
          "Success is expected to be handled on caller side");
   switch (Error) {
+  case PI_ERROR_OUT_OF_RESOURCES:
+    return handleOutOfResources(DeviceImpl, Kernel, NDRDesc);
+
   case PI_ERROR_INVALID_WORK_GROUP_SIZE:
     return handleInvalidWorkGroupSize(DeviceImpl, Kernel, NDRDesc);
 

sycl/test-e2e/OptionalKernelFeatures/throw-exception-for-out-of-registers-on-kernel-launch.cpp

@@ -26,7 +26,7 @@
 class kernel_vadd_and_sum;
 
 int main() {
-  sycl::queue q;
+  sycl::queue q{};
   sycl::device dev = q.get_device();
   size_t local_size = dev.get_info<sycl::info::device::max_work_group_size>();
   if (local_size < 1024u) {
@@ -80,6 +80,8 @@ int main() {
              // compute vector add
              const auto vadd = values1 + values2 + values3 + values4;
 
+             // NB: 64 registers used to do the vector addition.
+
              // compute total vector elements sum
              auto sum = elem_t(0);
              for (int j = 0; j < VEC_DIM; j++) {
@@ -92,11 +94,13 @@ int main() {
              output[i] = vadd;
              output[i] += sum;
            });
-     }).wait();
-  } catch (sycl::exception &e) {
+     }).wait_and_throw();
+  } catch (const sycl::exception &e) {
     using std::string_view_literals::operator""sv;
     auto Msg = "Exceeded the number of registers available on the hardware."sv;
-    if (std::string(e.what()).find(Msg) != std::string::npos) {
+    auto Errc = sycl::make_error_code(sycl::errc::nd_range);
+    if (e.code() == Errc &&
+        std::string_view{e.what()}.find(Msg) != std::string_view::npos) {
       return 0;
     }
   }