Skip to content

Commit 3fc0d49

Browse files
RossBruntonRossBrunton
committed
Update spec to make kernel validation optional
Several adapters don't support validating kernel signatures when enqueued. To handle this, we now allow urEnqueueKernelLaunch to return `SUCCESS` even when parameters are invalid. Some tests have also been updated. The CUDA adapter has also been updated to handle invalid arguments.
1 parent 59a1bab commit 3fc0d49

File tree

5 files changed

+40
-26
lines changed

5 files changed

+40
-26
lines changed

‎unified-runtime/include/ur_api.h

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -7431,6 +7431,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventSetCallback(
74317431
///////////////////////////////////////////////////////////////////////////////
74327432
/// @brief Enqueue a command to execute a kernel
74337433
///
7434+
/// @details
7435+
/// - Adapters may perform validation on the number of arguments set to the
7436+
/// kernel, but are not required to do so and may return
7437+
/// `::UR_RESULT_SUCCESS` even for invalid invocations.
7438+
///
74347439
/// @remarks
74357440
/// _Analogues_
74367441
/// - **clEnqueueNDRangeKernel**
@@ -7458,8 +7463,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventSetCallback(
74587463
/// - ::UR_RESULT_ERROR_INVALID_WORK_DIMENSION
74597464
/// - ::UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE
74607465
/// - ::UR_RESULT_ERROR_INVALID_VALUE
7461-
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS - "The kernel argument values
7462-
/// have not been specified."
7466+
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS
7467+
/// + The kernel argument values have not been specified and the adapter
7468+
/// is able to detect this.
74637469
/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY
74647470
/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES
74657471
UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(

‎unified-runtime/scripts/core/enqueue.yml

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,9 @@ type: function
1616
desc: "Enqueue a command to execute a kernel"
1717
class: $xEnqueue
1818
name: KernelLaunch
19+
details:
20+
- "Adapters may perform validation on the number of arguments set to the kernel, but are not required to do so and may
21+
return `$X_RESULT_SUCCESS` even for invalid invocations."
1922
ordinal: "0"
2023
analogue:
2124
- "**clEnqueueNDRangeKernel**"
@@ -65,8 +68,8 @@ returns:
6568
- $X_RESULT_ERROR_INVALID_WORK_DIMENSION
6669
- $X_RESULT_ERROR_INVALID_WORK_GROUP_SIZE
6770
- $X_RESULT_ERROR_INVALID_VALUE
68-
- $X_RESULT_ERROR_INVALID_KERNEL_ARGS
69-
- "The kernel argument values have not been specified."
71+
- $X_RESULT_ERROR_INVALID_KERNEL_ARGS:
72+
- "The kernel argument values have not been specified and the adapter is able to detect this."
7073
- $X_RESULT_ERROR_OUT_OF_HOST_MEMORY
7174
- $X_RESULT_ERROR_OUT_OF_RESOURCES
7275
--- #--------------------------------------------------------------------------

‎unified-runtime/source/loader/ur_libapi.cpp

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4982,6 +4982,11 @@ ur_result_t UR_APICALL urEventSetCallback(
49824982
///////////////////////////////////////////////////////////////////////////////
49834983
/// @brief Enqueue a command to execute a kernel
49844984
///
4985+
/// @details
4986+
/// - Adapters may perform validation on the number of arguments set to the
4987+
/// kernel, but are not required to do so and may return
4988+
/// `::UR_RESULT_SUCCESS` even for invalid invocations.
4989+
///
49854990
/// @remarks
49864991
/// _Analogues_
49874992
/// - **clEnqueueNDRangeKernel**
@@ -5009,8 +5014,9 @@ ur_result_t UR_APICALL urEventSetCallback(
50095014
/// - ::UR_RESULT_ERROR_INVALID_WORK_DIMENSION
50105015
/// - ::UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE
50115016
/// - ::UR_RESULT_ERROR_INVALID_VALUE
5012-
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS - "The kernel argument values
5013-
/// have not been specified."
5017+
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS
5018+
/// + The kernel argument values have not been specified and the adapter
5019+
/// is able to detect this.
50145020
/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY
50155021
/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES
50165022
ur_result_t UR_APICALL urEnqueueKernelLaunch(

‎unified-runtime/source/ur_api.cpp

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4347,6 +4347,11 @@ ur_result_t UR_APICALL urEventSetCallback(
43474347
///////////////////////////////////////////////////////////////////////////////
43484348
/// @brief Enqueue a command to execute a kernel
43494349
///
4350+
/// @details
4351+
/// - Adapters may perform validation on the number of arguments set to the
4352+
/// kernel, but are not required to do so and may return
4353+
/// `::UR_RESULT_SUCCESS` even for invalid invocations.
4354+
///
43504355
/// @remarks
43514356
/// _Analogues_
43524357
/// - **clEnqueueNDRangeKernel**
@@ -4374,8 +4379,9 @@ ur_result_t UR_APICALL urEventSetCallback(
43744379
/// - ::UR_RESULT_ERROR_INVALID_WORK_DIMENSION
43754380
/// - ::UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE
43764381
/// - ::UR_RESULT_ERROR_INVALID_VALUE
4377-
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS - "The kernel argument values
4378-
/// have not been specified."
4382+
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS
4383+
/// + The kernel argument values have not been specified and the adapter
4384+
/// is able to detect this.
43794385
/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY
43804386
/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES
43814387
ur_result_t UR_APICALL urEnqueueKernelLaunch(

‎unified-runtime/test/conformance/enqueue/urEnqueueKernelLaunch.cpp

Lines changed: 11 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -154,26 +154,19 @@ TEST_P(urEnqueueKernelLaunchTest, InvalidWorkGroupSize) {
154154
}
155155

156156
TEST_P(urEnqueueKernelLaunchTest, InvalidKernelArgs) {
157-
// Cuda and hip both lack any way to validate kernel args
158-
UUR_KNOWN_FAILURE_ON(uur::CUDA{}, uur::HIP{});
159-
UUR_KNOWN_FAILURE_ON(uur::LevelZero{}, uur::LevelZeroV2{});
160-
161-
ur_platform_backend_t backend;
162-
ASSERT_SUCCESS(urPlatformGetInfo(platform, UR_PLATFORM_INFO_BACKEND,
163-
sizeof(ur_platform_backend_t), &backend,
164-
nullptr));
165-
166-
if (backend == UR_PLATFORM_BACKEND_CUDA ||
167-
backend == UR_PLATFORM_BACKEND_HIP ||
168-
backend == UR_PLATFORM_BACKEND_LEVEL_ZERO) {
169-
GTEST_FAIL() << "AMD, L0 and Nvidia can't check kernel arguments.";
170-
}
157+
// Seems to segfault
158+
UUR_KNOWN_FAILURE_ON(uur::HIP{});
159+
// cuLaunchKernel seems to be returning CUDA_ERROR_INVALID_VALUE which is
160+
// converted to UR_RESULT_ERROR_INVALID_VALUE
161+
// https://github.com/oneapi-src/unified-runtime/issues/2720
162+
UUR_KNOWN_FAILURE_ON(uur::CUDA{});
171163

172164
// Enqueue kernel without setting any args
173-
ASSERT_EQ_RESULT(urEnqueueKernelLaunch(queue, kernel, n_dimensions,
174-
&global_offset, &global_size, nullptr,
175-
0, nullptr, nullptr),
176-
UR_RESULT_ERROR_INVALID_KERNEL_ARGS);
165+
auto error =
166+
urEnqueueKernelLaunch(queue, kernel, n_dimensions, &global_offset,
167+
&global_size, nullptr, 0, nullptr, nullptr);
168+
ASSERT_TRUE(error == UR_RESULT_ERROR_INVALID_KERNEL_ARGS ||
169+
error == UR_RESULT_SUCCESS);
177170
}
178171

179172
TEST_P(urEnqueueKernelLaunchKernelWgSizeTest, Success) {

0 commit comments

Comments
 (0)