Make invalid kernel arguments UB (#17962)

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

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

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
 --- #--------------------------------------------------------------------------

scripts/core/common.yml

@@ -86,7 +86,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:

scripts/core/enqueue.yml

@@ -65,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;

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:

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{});