Skip to content

Commit 9f6415d

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 5f48cdb commit 9f6415d

File tree

5 files changed

+39
-26
lines changed

5 files changed

+39
-26
lines changed

‎include/ur_api.h

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -7429,6 +7429,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventSetCallback(
74297429
///////////////////////////////////////////////////////////////////////////////
74307430
/// @brief Enqueue a command to execute a kernel
74317431
///
7432+
/// @details
7433+
/// - Adapters may perform validation on the number of arguments set to the
7434+
/// kernel, but are not required to do so and may return
7435+
/// `::UR_RESULT_SUCCESS` even for invalid invocations.
7436+
///
74327437
/// @remarks
74337438
/// _Analogues_
74347439
/// - **clEnqueueNDRangeKernel**
@@ -7456,8 +7461,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventSetCallback(
74567461
/// - ::UR_RESULT_ERROR_INVALID_WORK_DIMENSION
74577462
/// - ::UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE
74587463
/// - ::UR_RESULT_ERROR_INVALID_VALUE
7459-
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS - "The kernel argument values
7460-
/// have not been specified."
7464+
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS
7465+
/// + The kernel argument values have not been specified and the adapter
7466+
/// is able to detect this.
74617467
/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY
74627468
/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES
74637469
UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(

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

‎source/loader/ur_libapi.cpp

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4980,6 +4980,11 @@ ur_result_t UR_APICALL urEventSetCallback(
49804980
///////////////////////////////////////////////////////////////////////////////
49814981
/// @brief Enqueue a command to execute a kernel
49824982
///
4983+
/// @details
4984+
/// - Adapters may perform validation on the number of arguments set to the
4985+
/// kernel, but are not required to do so and may return
4986+
/// `::UR_RESULT_SUCCESS` even for invalid invocations.
4987+
///
49834988
/// @remarks
49844989
/// _Analogues_
49854990
/// - **clEnqueueNDRangeKernel**
@@ -5007,8 +5012,9 @@ ur_result_t UR_APICALL urEventSetCallback(
50075012
/// - ::UR_RESULT_ERROR_INVALID_WORK_DIMENSION
50085013
/// - ::UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE
50095014
/// - ::UR_RESULT_ERROR_INVALID_VALUE
5010-
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS - "The kernel argument values
5011-
/// have not been specified."
5015+
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS
5016+
/// + The kernel argument values have not been specified and the adapter
5017+
/// is able to detect this.
50125018
/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY
50135019
/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES
50145020
ur_result_t UR_APICALL urEnqueueKernelLaunch(

‎source/ur_api.cpp

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4345,6 +4345,11 @@ ur_result_t UR_APICALL urEventSetCallback(
43454345
///////////////////////////////////////////////////////////////////////////////
43464346
/// @brief Enqueue a command to execute a kernel
43474347
///
4348+
/// @details
4349+
/// - Adapters may perform validation on the number of arguments set to the
4350+
/// kernel, but are not required to do so and may return
4351+
/// `::UR_RESULT_SUCCESS` even for invalid invocations.
4352+
///
43484353
/// @remarks
43494354
/// _Analogues_
43504355
/// - **clEnqueueNDRangeKernel**
@@ -4372,8 +4377,9 @@ ur_result_t UR_APICALL urEventSetCallback(
43724377
/// - ::UR_RESULT_ERROR_INVALID_WORK_DIMENSION
43734378
/// - ::UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE
43744379
/// - ::UR_RESULT_ERROR_INVALID_VALUE
4375-
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS - "The kernel argument values
4376-
/// have not been specified."
4380+
/// - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGS
4381+
/// + The kernel argument values have not been specified and the adapter
4382+
/// is able to detect this.
43774383
/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY
43784384
/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES
43794385
ur_result_t UR_APICALL urEnqueueKernelLaunch(

‎test/conformance/enqueue/urEnqueueKernelLaunch.cpp

Lines changed: 10 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -154,26 +154,18 @@ 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+
UUR_KNOWN_FAILURE_ON(uur::CUDA{});
171162

172163
// 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);
164+
auto error =
165+
urEnqueueKernelLaunch(queue, kernel, n_dimensions, &global_offset,
166+
&global_size, nullptr, 0, nullptr, nullptr);
167+
ASSERT_TRUE(error == UR_RESULT_ERROR_INVALID_KERNEL_ARGS ||
168+
error == UR_RESULT_SUCCESS);
177169
}
178170

179171
TEST_P(urEnqueueKernelLaunchKernelWgSizeTest, Success) {

0 commit comments

Comments
 (0)