[UR] Add UR_DEVICE_INFO_VALIDATES_ON_LAUNCH

This device info is used to specify whether a given device can validate
arguments passed to a kernel when it is launched. This also updates the
spec to specify that if this value is not set, then running a kernel
with invalid arguments is UB.

This does not imply any behavior changes in adapters themselves; they
already UB-ed anyway. This just writes it into the spec and makes
some conformance tests not fail them for it.

Author: RossBrunton

unified-runtime/include/ur_api.h

@@ -462,6 +462,8 @@ etors:
       desc: "[int32_t][optional-query] return max power limit in milliwatts."
     - name: BFLOAT16_CONVERSIONS_NATIVE
       desc: "[$x_bool_t] support for native bfloat16 conversions"
+    - name: VALIDATES_ON_LAUNCH
+      desc: "[$x_bool_t] $xEnqueueKernelLaunch may return $X_RESULT_ERROR_INVALID_KERNEL_ARGS"
 --- #--------------------------------------------------------------------------
 type: function
 desc: "Retrieves various information about device"

unified-runtime/include/ur_print.hpp

@@ -16,6 +16,10 @@ type: function
 desc: "Enqueue a command to execute a kernel"
 class: $xEnqueue
 name: KernelLaunch
+details:
+    - "Adapters must perform validation on provided kernel arguments when launching them if they report
+      `$X_DEVICE_INFO_VALIDATES_ON_LAUNCH` as true for the appropriate device. If this is reported as false, then
+      launching a kernel with invalid arguments is Undefined Behavior."
 ordinal: "0"
 analogue:
     - "**clEnqueueNDRangeKernel**"
@@ -65,8 +69,8 @@ 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_INVALID_KERNEL_ARGS:
+        - "The kernel argument values are not valid and `$X_DEVICE_INFO_VALIDATE_ON_LAUNCH` is true for the device"
     - $X_RESULT_ERROR_OUT_OF_HOST_MEMORY
     - $X_RESULT_ERROR_OUT_OF_RESOURCES
 --- #--------------------------------------------------------------------------

unified-runtime/scripts/core/device.yml

@@ -1188,6 +1188,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice,
     return ReturnValue(false);
   case UR_DEVICE_INFO_ASYNC_USM_ALLOCATIONS_SUPPORT_EXP:
     return ReturnValue(true);
+  case UR_DEVICE_INFO_VALIDATES_ON_LAUNCH:
+    return ReturnValue(false);
   default:
     break;
   }

unified-runtime/scripts/core/enqueue.yml

@@ -1095,6 +1095,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice,
     return ReturnValue(true);
   case UR_DEVICE_INFO_MULTI_DEVICE_COMPILE_SUPPORT_EXP:
     return ReturnValue(false);
+    case UR_DEVICE_INFO_VALIDATES_ON_LAUNCH:
+      return ReturnValue(false);
   default:
     break;
   }

unified-runtime/source/adapters/cuda/device.cpp

@@ -1335,6 +1335,8 @@ ur_result_t urDeviceGetInfo(
       return ReturnValue(int32_t{PowerProperties.maxLimit});
     }
   }
+  case UR_DEVICE_INFO_VALIDATES_ON_LAUNCH:
+    return ReturnValue(false);
   default:
     logger::error("Unsupported ParamName in urGetDeviceInfo");
     logger::error("ParamNameParamName={}(0x{})", ParamName,

unified-runtime/source/adapters/hip/device.cpp

@@ -456,6 +456,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice,
   case UR_DEVICE_INFO_GLOBAL_VARIABLE_SUPPORT:
     return ReturnValue(false);
 
+  case UR_DEVICE_INFO_VALIDATES_ON_LAUNCH:
+    return ReturnValue(false);
+
   default:
     DIE_NO_IMPLEMENTATION;
   }

unified-runtime/source/adapters/level_zero/device.cpp

@@ -1427,6 +1427,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice,
     return ReturnValue(true);
   case UR_DEVICE_INFO_MULTI_DEVICE_COMPILE_SUPPORT_EXP:
     return ReturnValue(false);
+  case UR_DEVICE_INFO_VALIDATES_ON_LAUNCH:
+    return ReturnValue(true);
   // TODO: We can't query to check if these are supported, they will need to be
   // manually updated if support is ever implemented.
   case UR_DEVICE_INFO_KERNEL_SET_SPECIALIZATION_CONSTANTS:

unified-runtime/source/adapters/native_cpu/device.cpp

@@ -2662,6 +2662,20 @@ TEST_P(urDeviceGetInfoTest, SuccessMinPowerLimit) {
                              property_value);
 }
 
+TEST_P(urDeviceGetInfoTest, SuccessValidatesOnLaunch) {
+  size_t property_size = 0;
+  const ur_device_info_t property_name = UR_DEVICE_INFO_VALIDATES_ON_LAUNCH;
+
+  ASSERT_SUCCESS(
+      urDeviceGetInfo(device, property_name, 0, nullptr, &property_size));
+
+  ASSERT_EQ(property_size, sizeof(ur_bool_t));
+
+  ur_bool_t property_value = 0;
+  ASSERT_SUCCESS(urDeviceGetInfo(device, property_name, property_size,
+                                 &property_value, nullptr));
+}
+
 TEST_P(urDeviceGetInfoTest, InvalidNullHandleDevice) {
   ur_device_type_t device_type;
   ASSERT_EQ_RESULT(UR_RESULT_ERROR_INVALID_NULL_HANDLE,

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

@@ -154,19 +154,13 @@ TEST_P(urEnqueueKernelLaunchTest, InvalidWorkGroupSize) {
 }
 
 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));
+  ur_bool_t validates;
+  ASSERT_SUCCESS(urDeviceGetInfo(device, UR_DEVICE_INFO_VALIDATES_ON_LAUNCH,
+                                 sizeof(validates), &validates, 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.";
+  if (!validates) {
+    GTEST_SKIP() << "Adapter can't check kernel arguments.";
+    return;
   }
 
   // Enqueue kernel without setting any args