[UR] Make invalid kernel arguments UB

Instead of returning INVALID_KERNEL_ARGS, calling a kernel with invalid
arguments is now UB. In addition, the Invalid kernel arguments enum
itself is now deprecated (as it is no longer returned) and the OpenCL
backend now writes a message to the log.

Author: RossBrunton

sycl/source/detail/error_handling/error_handling.cpp

@@ -404,14 +404,6 @@ void handleErrorOrWarning(ur_result_t Error, const device_impl &DeviceImpl,
   case UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE:
     return handleInvalidWorkGroupSize(DeviceImpl, Kernel, NDRDesc);
 
-  case UR_RESULT_ERROR_INVALID_KERNEL_ARGS:
-    throw detail::set_ur_error(
-        sycl::exception(
-            make_error_code(errc::kernel_argument),
-            "The kernel argument values have not been specified OR a kernel "
-            "argument declared to be a pointer to a type."),
-        UR_RESULT_ERROR_INVALID_KERNEL_ARGS);
-
   case UR_RESULT_ERROR_INVALID_WORK_ITEM_SIZE:
     return handleInvalidWorkItemSize(DeviceImpl, NDRDesc);
 

sycl/source/exception.cpp

@@ -103,7 +103,6 @@ __SYCL_EXPORT const char *stringifyErrorCode(int32_t error) {
     _UR_ERRC(UR_RESULT_ERROR_INVALID_DEVICE_PARTITION_COUNT)
     _UR_ERRC(UR_RESULT_ERROR_INVALID_WORK_ITEM_SIZE)
     _UR_ERRC(UR_RESULT_ERROR_INVALID_WORK_DIMENSION)
-    _UR_ERRC(UR_RESULT_ERROR_INVALID_KERNEL_ARGS)
     _UR_ERRC(UR_RESULT_ERROR_INVALID_KERNEL)
     _UR_ERRC(UR_RESULT_ERROR_INVALID_KERNEL_NAME)
     _UR_ERRC(UR_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX)

unified-runtime/include/ur_api.h

@@ -196,7 +196,7 @@ etors:
     - name: ERROR_INVALID_WORK_DIMENSION
       desc: "Invalid work dimension"
     - name: ERROR_INVALID_KERNEL_ARGS
-      desc: "Invalid kernel args"
+      desc: "[deprecated-value] No longer used - invalid kernel args are now UB"
     - name: ERROR_INVALID_KERNEL
       desc: "Invalid kernel"
     - name: ERROR_INVALID_KERNEL_NAME

unified-runtime/include/ur_print.hpp

@@ -19,6 +19,8 @@ name: KernelLaunch
 ordinal: "0"
 analogue:
     - "**clEnqueueNDRangeKernel**"
+details:
+    - "Providing invalid kernel arguments is Undefined Behavior."
 params:
     - type: $x_queue_handle_t
       name: hQueue
@@ -65,8 +67,6 @@ returns:
     - $X_RESULT_ERROR_INVALID_WORK_DIMENSION
     - $X_RESULT_ERROR_INVALID_WORK_GROUP_SIZE
     - $X_RESULT_ERROR_INVALID_VALUE
-    - $X_RESULT_ERROR_INVALID_KERNEL_ARGS
-        - "The kernel argument values have not been specified."
     - $X_RESULT_ERROR_OUT_OF_HOST_MEMORY
     - $X_RESULT_ERROR_OUT_OF_RESOURCES
 --- #--------------------------------------------------------------------------

unified-runtime/scripts/core/common.yml

@@ -94,7 +94,7 @@ ur_result_t mapCLErrorToUR(cl_int Result) {
   case CL_DEVICE_NOT_AVAILABLE:
     return UR_RESULT_ERROR_DEVICE_NOT_AVAILABLE;
   case CL_INVALID_KERNEL_ARGS:
-    return UR_RESULT_ERROR_INVALID_KERNEL_ARGS;
+    return UR_RESULT_ERROR_ADAPTER_SPECIFIC;
   case CL_INVALID_COMMAND_QUEUE:
     return UR_RESULT_ERROR_INVALID_QUEUE;
   case CL_INVALID_ARG_SIZE:

unified-runtime/scripts/core/enqueue.yml

@@ -8,6 +8,7 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include "adapter.hpp"
 #include "common.hpp"
 #include "context.hpp"
 #include "event.hpp"
@@ -64,12 +65,16 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
   cl_event Event;
   std::vector<cl_event> CLWaitEvents(numEventsInWaitList);
   MapUREventsToCL(numEventsInWaitList, phEventWaitList, CLWaitEvents);
-  CL_RETURN_ON_FAILURE(clEnqueueNDRangeKernel(
+  auto Err = clEnqueueNDRangeKernel(
       hQueue->CLQueue, hKernel->CLKernel, workDim, pGlobalWorkOffset,
       pGlobalWorkSize,
       compiledLocalWorksize.empty() ? pLocalWorkSize
                                     : compiledLocalWorksize.data(),
-      numEventsInWaitList, CLWaitEvents.data(), ifUrEvent(phEvent, Event)));
+      numEventsInWaitList, CLWaitEvents.data(), ifUrEvent(phEvent, Event));
+  if (Err == CL_INVALID_KERNEL_ARGS) {
+    ur::cl::getAdapter()->log.error("Kernel called with invalid arguments");
+  }
+  CL_RETURN_ON_FAILURE(Err);
 
   UR_RETURN_ON_FAILURE(createUREvent(Event, hQueue->Context, hQueue, phEvent));
   return UR_RESULT_SUCCESS;

unified-runtime/source/adapters/opencl/common.cpp

@@ -123,8 +123,6 @@ static auto getUrResultString = [](ur_result_t Result) {
     return "UR_RESULT_ERROR_INVALID_WORK_ITEM_SIZE";
   case UR_RESULT_ERROR_INVALID_WORK_DIMENSION:
     return "UR_RESULT_ERROR_INVALID_WORK_DIMENSION";
-  case UR_RESULT_ERROR_INVALID_KERNEL_ARGS:
-    return "UR_RESULT_ERROR_INVALID_KERNEL_ARGS";
   case UR_RESULT_ERROR_INVALID_KERNEL:
     return "UR_RESULT_ERROR_INVALID_KERNEL";
   case UR_RESULT_ERROR_INVALID_KERNEL_NAME:

unified-runtime/source/adapters/opencl/enqueue.cpp

@@ -153,29 +153,6 @@ TEST_P(urEnqueueKernelLaunchTest, InvalidWorkGroupSize) {
               result == UR_RESULT_SUCCESS);
 }
 
-TEST_P(urEnqueueKernelLaunchTest, InvalidKernelArgs) {
-  // Cuda and hip both lack any way to validate kernel args
-  UUR_KNOWN_FAILURE_ON(uur::CUDA{}, uur::HIP{});
-  UUR_KNOWN_FAILURE_ON(uur::LevelZero{}, uur::LevelZeroV2{});
-
-  ur_platform_backend_t backend;
-  ASSERT_SUCCESS(urPlatformGetInfo(platform, UR_PLATFORM_INFO_BACKEND,
-                                   sizeof(ur_platform_backend_t), &backend,
-                                   nullptr));
-
-  if (backend == UR_PLATFORM_BACKEND_CUDA ||
-      backend == UR_PLATFORM_BACKEND_HIP ||
-      backend == UR_PLATFORM_BACKEND_LEVEL_ZERO) {
-    GTEST_FAIL() << "AMD, L0 and Nvidia can't check kernel arguments.";
-  }
-
-  // Enqueue kernel without setting any args
-  ASSERT_EQ_RESULT(urEnqueueKernelLaunch(queue, kernel, n_dimensions,
-                                         &global_offset, &global_size, nullptr,
-                                         0, nullptr, nullptr),
-                   UR_RESULT_ERROR_INVALID_KERNEL_ARGS);
-}
-
 TEST_P(urEnqueueKernelLaunchKernelWgSizeTest, Success) {
   UUR_KNOWN_FAILURE_ON(uur::LevelZero{}, uur::LevelZeroV2{});