diff --git a/unified-runtime/include/ur_api.h b/unified-runtime/include/ur_api.h index c390ed4410d16..544f3678d28fc 100644 --- a/unified-runtime/include/ur_api.h +++ b/unified-runtime/include/ur_api.h @@ -7431,6 +7431,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventSetCallback( /////////////////////////////////////////////////////////////////////////////// /// @brief Enqueue a command to execute a kernel /// +/// @details +/// - Adapters may perform validation on the number of arguments set to the +/// kernel, but are not required to do so and may return +/// `::UR_RESULT_SUCCESS` even for invalid invocations. +/// /// @remarks /// _Analogues_ /// - **clEnqueueNDRangeKernel** @@ -7458,8 +7463,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventSetCallback( /// - ::UR_RESULT_ERROR_INVALID_WORK_DIMENSION /// - ::UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE /// - ::UR_RESULT_ERROR_INVALID_VALUE -/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS - "The kernel argument values -/// have not been specified." +/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS +/// + The kernel argument values have not been specified and the adapter +/// is able to detect this. /// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY /// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch( diff --git a/unified-runtime/scripts/core/enqueue.yml b/unified-runtime/scripts/core/enqueue.yml index cc2597962d338..36add449bf8e5 100644 --- a/unified-runtime/scripts/core/enqueue.yml +++ b/unified-runtime/scripts/core/enqueue.yml @@ -16,6 +16,9 @@ type: function desc: "Enqueue a command to execute a kernel" class: $xEnqueue name: KernelLaunch +details: + - "Adapters may perform validation on the number of arguments set to the kernel, but are not required to do so and may + return `$X_RESULT_SUCCESS` even for invalid invocations." ordinal: "0" analogue: - "**clEnqueueNDRangeKernel**" @@ -65,8 +68,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 have not been specified and the adapter is able to detect this." - $X_RESULT_ERROR_OUT_OF_HOST_MEMORY - $X_RESULT_ERROR_OUT_OF_RESOURCES --- #-------------------------------------------------------------------------- diff --git a/unified-runtime/source/loader/ur_libapi.cpp b/unified-runtime/source/loader/ur_libapi.cpp index e5797537632bf..5e8d0bbd2cc61 100644 --- a/unified-runtime/source/loader/ur_libapi.cpp +++ b/unified-runtime/source/loader/ur_libapi.cpp @@ -4982,6 +4982,11 @@ ur_result_t UR_APICALL urEventSetCallback( /////////////////////////////////////////////////////////////////////////////// /// @brief Enqueue a command to execute a kernel /// +/// @details +/// - Adapters may perform validation on the number of arguments set to the +/// kernel, but are not required to do so and may return +/// `::UR_RESULT_SUCCESS` even for invalid invocations. +/// /// @remarks /// _Analogues_ /// - **clEnqueueNDRangeKernel** @@ -5009,8 +5014,9 @@ ur_result_t UR_APICALL urEventSetCallback( /// - ::UR_RESULT_ERROR_INVALID_WORK_DIMENSION /// - ::UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE /// - ::UR_RESULT_ERROR_INVALID_VALUE -/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS - "The kernel argument values -/// have not been specified." +/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS +/// + The kernel argument values have not been specified and the adapter +/// is able to detect this. /// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY /// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES ur_result_t UR_APICALL urEnqueueKernelLaunch( diff --git a/unified-runtime/source/ur_api.cpp b/unified-runtime/source/ur_api.cpp index c5651c0fc4f83..f7f554e7a7d17 100644 --- a/unified-runtime/source/ur_api.cpp +++ b/unified-runtime/source/ur_api.cpp @@ -4347,6 +4347,11 @@ ur_result_t UR_APICALL urEventSetCallback( /////////////////////////////////////////////////////////////////////////////// /// @brief Enqueue a command to execute a kernel /// +/// @details +/// - Adapters may perform validation on the number of arguments set to the +/// kernel, but are not required to do so and may return +/// `::UR_RESULT_SUCCESS` even for invalid invocations. +/// /// @remarks /// _Analogues_ /// - **clEnqueueNDRangeKernel** @@ -4374,8 +4379,9 @@ ur_result_t UR_APICALL urEventSetCallback( /// - ::UR_RESULT_ERROR_INVALID_WORK_DIMENSION /// - ::UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE /// - ::UR_RESULT_ERROR_INVALID_VALUE -/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS - "The kernel argument values -/// have not been specified." +/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS +/// + The kernel argument values have not been specified and the adapter +/// is able to detect this. /// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY /// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES ur_result_t UR_APICALL urEnqueueKernelLaunch( diff --git a/unified-runtime/test/conformance/enqueue/urEnqueueKernelLaunch.cpp b/unified-runtime/test/conformance/enqueue/urEnqueueKernelLaunch.cpp index ef5c0228ede48..173caae0bd1f3 100644 --- a/unified-runtime/test/conformance/enqueue/urEnqueueKernelLaunch.cpp +++ b/unified-runtime/test/conformance/enqueue/urEnqueueKernelLaunch.cpp @@ -154,26 +154,19 @@ 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)); - - 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."; - } + // Seems to segfault + UUR_KNOWN_FAILURE_ON(uur::HIP{}); + // cuLaunchKernel seems to be returning CUDA_ERROR_INVALID_VALUE which is + // converted to UR_RESULT_ERROR_INVALID_VALUE + // https://github.com/oneapi-src/unified-runtime/issues/2720 + UUR_KNOWN_FAILURE_ON(uur::CUDA{}); // 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); + auto error = + urEnqueueKernelLaunch(queue, kernel, n_dimensions, &global_offset, + &global_size, nullptr, 0, nullptr, nullptr); + ASSERT_TRUE(error == UR_RESULT_ERROR_INVALID_KERNEL_ARGS || + error == UR_RESULT_SUCCESS); } TEST_P(urEnqueueKernelLaunchKernelWgSizeTest, Success) {