From 34e9baecdae439414c06c589bf8189d025ad7027 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Thu, 27 Nov 2025 00:19:59 +0000 Subject: [PATCH] Test mutable command-buffer deferred kernel arguments (#2538) Tests the scenario outlined in https://github.com/KhronosGroup/OpenCL-CTS/issues/2520#issuecomment-3280930492 to verify the functionality added in OpenCL-Docs PR https://github.com/KhronosGroup/OpenCL-Docs/pull/1382. Closes #2520 --------- Co-authored-by: Ben Ashbaugh --- .../CMakeLists.txt | 1 + .../mutable_command_basic.h | 4 +- .../mutable_command_defer_arguments.cpp | 219 ++++++++++++++++++ 3 files changed, 222 insertions(+), 2 deletions(-) create mode 100644 test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_defer_arguments.cpp diff --git a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/CMakeLists.txt b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/CMakeLists.txt index aed183ff..bb085035 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/CMakeLists.txt +++ b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/CMakeLists.txt @@ -16,6 +16,7 @@ set(${MODULE_NAME}_SOURCES mutable_command_work_groups.cpp mutable_command_work_dim.cpp mutable_command_update_state.cpp + mutable_command_defer_arguments.cpp ../basic_command_buffer.cpp ) diff --git a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_basic.h b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_basic.h index 47b0c97a..917183c8 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_basic.h +++ b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_basic.h @@ -89,10 +89,10 @@ struct BasicMutableCommandBufferTest : BasicCommandBufferTest cl_version extension_version = get_extension_version( device, "cl_khr_command_buffer_mutable_dispatch"); - if (extension_version != CL_MAKE_VERSION(0, 9, 4)) + if (extension_version != CL_MAKE_VERSION(0, 9, 5)) { log_info("cl_khr_command_buffer_mutable_dispatch version " - "0.9.4 is " + "0.9.5 is " "required to run the test, skipping.\n "); extension_available = false; } diff --git a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_defer_arguments.cpp b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_defer_arguments.cpp new file mode 100644 index 00000000..2e10cbf8 --- /dev/null +++ b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_defer_arguments.cpp @@ -0,0 +1,219 @@ +// +// Copyright (c) 2025 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// + +#include "testHarness.h" +#include "mutable_command_basic.h" + +namespace { + +//////////////////////////////////////////////////////////////////////////////// +// Mutable dispatch test which handles the case where all the arguments of a +// kernel aren't set when a kernel is initially added to a mutable +// command-buffer, but deferred until an update is made to the command to set +// them before command-buffer enqueue. +struct MutableDispatchDeferArguments : public BasicMutableCommandBufferTest +{ + MutableDispatchDeferArguments(cl_device_id device, cl_context context, + cl_command_queue queue) + : BasicMutableCommandBufferTest(device, context, queue) + {} + + bool Skip() override + { + if (BasicMutableCommandBufferTest::Skip()) return true; + cl_mutable_dispatch_fields_khr mutable_capabilities; + bool mutable_support = + !clGetDeviceInfo( + device, CL_DEVICE_MUTABLE_DISPATCH_CAPABILITIES_KHR, + sizeof(mutable_capabilities), &mutable_capabilities, nullptr) + && mutable_capabilities & CL_MUTABLE_DISPATCH_ARGUMENTS_KHR; + + // require mutable arguments capability + return !mutable_support; + } + + cl_int SetUpKernel() override + { + // Create kernel + const char *defer_args_kernel = + R"( + __kernel void defer_args_test(__constant int *src, __global int *dst) + { + size_t tid = get_global_id(0); + dst[tid] = src[tid]; + })"; + + cl_int error = + create_single_kernel_helper(context, &program, &kernel, 1, + &defer_args_kernel, "defer_args_test"); + test_error(error, "Creating kernel failed"); + return CL_SUCCESS; + } + + cl_int SetUpKernelArgs() override + { + // Create and initialize buffers + MTdataHolder d(gRandomSeed); + + src_data.resize(num_elements); + for (size_t i = 0; i < num_elements; i++) + src_data[i] = (cl_int)genrand_int32(d); + + cl_int error = CL_SUCCESS; + in_mem = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, + num_elements * sizeof(cl_int), src_data.data(), + &error); + test_error(error, "Creating src buffer"); + + out_mem = clCreateBuffer(context, CL_MEM_READ_WRITE, + num_elements * sizeof(cl_int), NULL, &error); + test_error(error, "Creating initial dst buffer failed"); + + // Only set a single kernel argument, leaving argument at index 1 unset + error = clSetKernelArg(kernel, 0, sizeof(in_mem), &in_mem); + test_error(error, "Unable to set src kernel arguments"); + + return CL_SUCCESS; + } + + bool verify_state(cl_command_buffer_state_khr expected) + { + cl_command_buffer_state_khr state = ~cl_command_buffer_state_khr(0); + cl_int error = clGetCommandBufferInfoKHR( + command_buffer, CL_COMMAND_BUFFER_STATE_KHR, sizeof(state), &state, + nullptr); + if (error != CL_SUCCESS) + { + log_error("clGetCommandBufferInfoKHR failed: %d", error); + return false; + } + + if (state != expected) + { + log_error("Unexpected result of CL_COMMAND_BUFFER_STATE_KHR query. " + "Expected %u, but was %u\n", + expected, state); + return false; + } + return true; + } + + bool verify_result(const cl_mem &buffer) + { + std::vector data(num_elements); + cl_int error = + clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, data_size(), + data.data(), 0, nullptr, nullptr); + test_error(error, "clEnqueueReadBuffer failed"); + + for (size_t i = 0; i < num_elements; i++) + { + if (data[i] != src_data[i]) + { + log_error("Modified verification failed at index %zu: Got %d, " + "wanted %d\n", + i, data[i], src_data[i]); + return false; + } + } + return true; + } + + cl_int Run() override + { + // Create command while the kernel still has the second argument unset. + // Passing 'CL_MUTABLE_DISPATCH_ARGUMENTS_KHR' as a property means this + // shouldn't be an error. + cl_command_properties_khr props[] = { + CL_MUTABLE_DISPATCH_UPDATABLE_FIELDS_KHR, + CL_MUTABLE_DISPATCH_ARGUMENTS_KHR, 0 + }; + cl_int error = clCommandNDRangeKernelKHR( + command_buffer, nullptr, props, kernel, 1, nullptr, &num_elements, + nullptr, 0, nullptr, nullptr, &command); + test_error(error, "clCommandNDRangeKernelKHR failed"); + + // Finalizing the command buffer shouldn't be an error, but result in + // the command-buffer entering the CL_COMMAND_BUFFER_STATE_FINALIZED + // state. + error = clFinalizeCommandBufferKHR(command_buffer); + test_error(error, "clFinalizeCommandBufferKHR failed"); + if (!verify_state(CL_COMMAND_BUFFER_STATE_FINALIZED_KHR)) + { + return TEST_FAIL; + } + + // Check that trying to enqueue the command-buffer in this state is an + // error, as it needs to be in the executable state. + error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0, + nullptr, nullptr); + test_failure_error_ret(error, CL_INVALID_OPERATION, + "clEnqueueCommandBufferKHR should return " + "CL_INVALID_OPERATION", + TEST_FAIL); + + // Update the kernel command to set the missing argument. + cl_mutable_dispatch_arg_khr arg{ 1, sizeof(out_mem), &out_mem }; + cl_mutable_dispatch_config_khr dispatch_config{ + command, + 1 /* num_args */, + 0 /* num_svm_arg */, + 0 /* num_exec_infos */, + 0 /* work_dim - 0 means no change to dimensions */, + &arg /* arg_list */, + nullptr /* arg_svm_list - nullptr means no change*/, + nullptr /* exec_info_list */, + nullptr /* global_work_offset */, + nullptr /* global_work_size */, + nullptr /* local_work_size */ + }; + + cl_uint num_configs = 1; + cl_command_buffer_update_type_khr config_types[1] = { + CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR + }; + const void *configs[1] = { &dispatch_config }; + error = clUpdateMutableCommandsKHR(command_buffer, num_configs, + config_types, configs); + test_error(error, "clUpdateMutableCommandsKHR failed"); + + // Now that all the arguments have been set, verify the + // command-buffer has entered the executable state. + if (!verify_state(CL_COMMAND_BUFFER_STATE_EXECUTABLE_KHR)) + { + return TEST_FAIL; + } + + // Execute command-buffer and verify results are expected + error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0, + nullptr, nullptr); + test_error(error, "clEnqueueCommandBufferKHR failed"); + if (!verify_result(out_mem)) return TEST_FAIL; + + return TEST_PASS; + } + + cl_mutable_command_khr command; + std::vector src_data; +}; + +} // anonymous namespace + +REGISTER_TEST(mutable_dispatch_defer_arguments) +{ + return MakeAndRunTest(device, context, queue, + num_elements); +}