diff --git a/sycl/test-e2e/Graph/Explicit/regression_accessor_spv.cpp b/sycl/test-e2e/Graph/Explicit/regression_accessor_spv.cpp index 56f19fe71e63a..ddb5d9604af6f 100644 --- a/sycl/test-e2e/Graph/Explicit/regression_accessor_spv.cpp +++ b/sycl/test-e2e/Graph/Explicit/regression_accessor_spv.cpp @@ -41,8 +41,10 @@ int main(int, char **argv) { Graph.add([&](handler &cgh) { auto Acc = BufA.get_access(cgh); - + // the kernel requires two arguments, the second one is an offset which is + // set to 0 here cgh.set_arg(0, Acc); + cgh.set_arg(1, 0); // offset cgh.single_task(kernel); }); diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_spv.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_spv.cpp index 6de2be2815ba7..5c889dfb69b1d 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_spv.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_accessor_spv.cpp @@ -38,6 +38,7 @@ int main(int, char **argv) { int PatternA = 42; int PatternB = 0xA; + int AccOffset = 0; auto AccA = BufA.get_access(); auto AccB = BufB.get_access(); @@ -52,6 +53,7 @@ int main(int, char **argv) { auto CGFA = [&](handler &CGH) { CGH.require(AccA); CGH.set_arg(0, AccA); + CGH.set_arg(1, AccOffset); CGH.set_arg(2, PatternA); CGH.parallel_for(sycl::range<1>(Size), kernelA); }; @@ -59,6 +61,7 @@ int main(int, char **argv) { auto CGFB = [&](handler &CGH) { CGH.require(AccB); CGH.set_arg(0, AccB); + CGH.set_arg(1, AccOffset); CGH.set_arg(2, PatternB); CGH.parallel_for(sycl::range<1>(Size), kernelB); }; diff --git a/sycl/test-e2e/Graph/Update/update_with_indices_accessor_spv.cpp b/sycl/test-e2e/Graph/Update/update_with_indices_accessor_spv.cpp index 2aaaa43a9ab73..1a188ef80c037 100644 --- a/sycl/test-e2e/Graph/Update/update_with_indices_accessor_spv.cpp +++ b/sycl/test-e2e/Graph/Update/update_with_indices_accessor_spv.cpp @@ -41,7 +41,10 @@ int main(int, char **argv) { auto KernelNode = Graph.add([&](handler &cgh) { cgh.require(InputParam); + // the kernel requires two arguments, the second one is an offset which is + // set to 0 here cgh.set_arg(0, InputParam); + cgh.set_arg(1, 0); // offset cgh.single_task(kernel); }); diff --git a/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp b/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp index 15fac6ee42d7d..127ebdb6faf20 100644 --- a/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp +++ b/unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp @@ -1172,6 +1172,13 @@ ur_result_t ur_command_list_manager::appendKernelLaunchWithArgsExpNew( const ur_kernel_launch_ext_properties_t *launchPropList, wait_list_view &waitListView, ur_event_handle_t phEvent) { + if (numArgs != hKernel->getCommonProperties().numKernelArgs) { + setErrorMessage("Wrong number of kernel arguments", + UR_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX, + static_cast(ZE_RESULT_ERROR_INVALID_ARGUMENT)); + return UR_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX; + } + ur_result_t checkResult = kernelLaunchChecks(hKernel, workDim); if (checkResult != UR_RESULT_SUCCESS) { return checkResult; @@ -1219,6 +1226,13 @@ ur_result_t ur_command_list_manager::appendKernelLaunchWithArgsExpNew( hKernel->kernelArgs.resize(numArgs, 0); for (uint32_t argIndex = 0; argIndex < numArgs; argIndex++) { + if (pArgs[argIndex].index != argIndex) { + setErrorMessage("Missing kernel argument", + UR_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX, + static_cast(ZE_RESULT_ERROR_INVALID_ARGUMENT)); + return UR_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX; + } + switch (pArgs[argIndex].type) { case UR_EXP_KERNEL_ARG_TYPE_LOCAL: hKernel->kernelArgs[argIndex] =