From 109c0a1ddd91feb0ff3653d53405f38c9b671d0a Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Tue, 19 Mar 2024 07:10:19 -0700 Subject: [PATCH 1/3] configure depdendabot github actions version updates (#1923) --- .github/dependabot.yml | 10 ++++++++++ 1 file changed, 10 insertions(+) create mode 100644 .github/dependabot.yml diff --git a/.github/dependabot.yml b/.github/dependabot.yml new file mode 100644 index 0000000000..2390d8c809 --- /dev/null +++ b/.github/dependabot.yml @@ -0,0 +1,10 @@ +version: 2 +updates: + - package-ecosystem: "github-actions" + directory: "/" + schedule: + interval: "monthly" + groups: + github-actions: + patterns: + - "*" From 8f3ef0891d51c89beecb804724e450b77c8e30ad Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Tue, 19 Mar 2024 16:41:41 +0100 Subject: [PATCH 2/3] Added new tests for simultaneous use with mutable dispatch (#1912) * Added new tests for simultaneous use with mutable dispatch -cross queue simultaneous use -in-order queue with simultaneous use According to issue description #1481 * Several corrections applied: -reordered Skip conditions to check valid simultaneous_use_support flag -removed unnecessary SetUpKernel call -initialize kernel and memory buffers from BasicCommandBufferTest instead BasicMutableCommandBufferTest * Corrections for command buffer creation to request simultaneous property --- .../CMakeLists.txt | 2 +- .../main.cpp | 2 + .../mutable_command_basic.h | 8 +- ...r.cpp => mutable_command_simultaneous.cpp} | 367 +++++++++++++----- .../procs.h | 7 + 5 files changed, 294 insertions(+), 92 deletions(-) rename test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/{mutable_command_out_of_order.cpp => mutable_command_simultaneous.cpp} (51%) 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 ecfe36f815..16f847d0e7 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 @@ -5,7 +5,7 @@ set(${MODULE_NAME}_SOURCES mutable_command_info.cpp mutable_command_image_arguments.cpp mutable_command_arguments.cpp - mutable_command_out_of_order.cpp + mutable_command_simultaneous.cpp mutable_command_global_size.cpp mutable_command_local_size.cpp mutable_command_global_offset.cpp diff --git a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/main.cpp b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/main.cpp index dbbdf8df45..e6af2898c9 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/main.cpp +++ b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/main.cpp @@ -31,6 +31,8 @@ test_definition test_list[] = { ADD_TEST(mutable_dispatch_image_2d_arguments), ADD_TEST(mutable_dispatch_out_of_order), ADD_TEST(mutable_dispatch_simultaneous_out_of_order), + ADD_TEST(mutable_dispatch_simultaneous_in_order), + ADD_TEST(mutable_dispatch_simultaneous_cross_queue), ADD_TEST(mutable_dispatch_global_size), ADD_TEST(mutable_dispatch_local_size), ADD_TEST(mutable_dispatch_global_offset), 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 a62e84b3ef..eee6a76ece 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 @@ -55,9 +55,15 @@ struct BasicMutableCommandBufferTest : BasicCommandBufferTest cl_int error = init_extension_functions(); test_error(error, "Unable to initialise extension functions"); + cl_command_buffer_properties_khr prop = CL_COMMAND_BUFFER_MUTABLE_KHR; + if (simultaneous_use_support) + { + prop |= CL_COMMAND_BUFFER_SIMULTANEOUS_USE_KHR; + } + const cl_command_buffer_properties_khr props[] = { CL_COMMAND_BUFFER_FLAGS_KHR, - CL_COMMAND_BUFFER_MUTABLE_KHR, + prop, 0, }; diff --git a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_out_of_order.cpp b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_simultaneous.cpp similarity index 51% rename from test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_out_of_order.cpp rename to test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_simultaneous.cpp index d507dadfa5..42dd90c74a 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_out_of_order.cpp +++ b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_simultaneous.cpp @@ -1,5 +1,5 @@ // -// Copyright (c) 2022 The Khronos Group Inc. +// Copyright (c) 2024 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. @@ -22,30 +22,30 @@ #include //////////////////////////////////////////////////////////////////////////////// // mutable dispatch tests which handle following cases: -// - simultaneous use -// - cross-queue simultaneous-use +// - out-of-order queue use +// - out-of-order queue with simultaneous use +// - in-order queue with simultaneous use +// - cross-queue with simultaneous use namespace { -template -struct OutOfOrderTest : public BasicMutableCommandBufferTest +template +struct SimultaneousMutableDispatchTest : public BasicMutableCommandBufferTest { - OutOfOrderTest(cl_device_id device, cl_context context, - cl_command_queue queue) + SimultaneousMutableDispatchTest(cl_device_id device, cl_context context, + cl_command_queue queue) : BasicMutableCommandBufferTest(device, context, queue), - out_of_order_queue(nullptr), out_of_order_command_buffer(this), - user_event(nullptr), wait_pass_event(nullptr), kernel_fill(nullptr), - program_fill(nullptr) + work_queue(nullptr), work_command_buffer(this), user_event(nullptr), + wait_pass_event(nullptr), command(nullptr) { simultaneous_use_requested = simultaneous_request; if (simultaneous_request) buffer_size_multiplier = 2; } - //-------------------------------------------------------------------------- cl_int SetUpKernel() override { - cl_int error = BasicMutableCommandBufferTest::SetUpKernel(); - test_error(error, "BasicMutableCommandBufferTest::SetUpKernel failed"); + cl_int error = BasicCommandBufferTest::SetUpKernel(); + test_error(error, "BasicCommandBufferTest::SetUpKernel failed"); // create additional kernel to properly prepare output buffer for test const char* kernel_str = @@ -72,12 +72,10 @@ struct OutOfOrderTest : public BasicMutableCommandBufferTest return CL_SUCCESS; } - //-------------------------------------------------------------------------- cl_int SetUpKernelArgs() override { - cl_int error = BasicMutableCommandBufferTest::SetUpKernelArgs(); - test_error(error, - "BasicMutableCommandBufferTest::SetUpKernelArgs failed"); + cl_int error = BasicCommandBufferTest::SetUpKernelArgs(); + test_error(error, "BasicCommandBufferTest::SetUpKernelArgs failed"); error = clSetKernelArg(kernel_fill, 0, sizeof(cl_int), &overwritten_pattern); @@ -92,33 +90,48 @@ struct OutOfOrderTest : public BasicMutableCommandBufferTest return CL_SUCCESS; } - //-------------------------------------------------------------------------- cl_int SetUp(int elements) override { cl_int error = BasicMutableCommandBufferTest::SetUp(elements); test_error(error, "BasicMutableCommandBufferTest::SetUp failed"); - error = SetUpKernel(); - test_error(error, "SetUpKernel failed"); + if (out_of_order_request) + { + work_queue = clCreateCommandQueue( + context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, + &error); + test_error(error, "Unable to create command queue to test with"); + + cl_command_buffer_properties_khr prop = + CL_COMMAND_BUFFER_MUTABLE_KHR; + if (simultaneous_use_support) + { + prop |= CL_COMMAND_BUFFER_SIMULTANEOUS_USE_KHR; + } - out_of_order_queue = clCreateCommandQueue( - context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &error); - test_error(error, "Unable to create command queue to test with"); + const cl_command_buffer_properties_khr props[] = { + CL_COMMAND_BUFFER_FLAGS_KHR, + prop, + 0, + }; - cl_command_buffer_properties_khr properties[3] = { - CL_COMMAND_BUFFER_FLAGS_KHR, CL_COMMAND_BUFFER_MUTABLE_KHR, 0 - }; - - out_of_order_command_buffer = clCreateCommandBufferKHR( - 1, &out_of_order_queue, properties, &error); - test_error(error, "clCreateCommandBufferKHR failed"); + work_command_buffer = + clCreateCommandBufferKHR(1, &work_queue, props, &error); + test_error(error, "clCreateCommandBufferKHR failed"); + } + else + { + work_queue = queue; + work_command_buffer = command_buffer; + } return CL_SUCCESS; } - //-------------------------------------------------------------------------- bool Skip() override { + if (BasicMutableCommandBufferTest::Skip()) return true; + cl_mutable_dispatch_fields_khr mutable_capabilities; bool mutable_support = @@ -127,13 +140,11 @@ struct OutOfOrderTest : public BasicMutableCommandBufferTest sizeof(mutable_capabilities), &mutable_capabilities, nullptr) && mutable_capabilities & CL_MUTABLE_DISPATCH_ARGUMENTS_KHR; - - return !out_of_order_support + return (out_of_order_request && !out_of_order_support) || (simultaneous_use_requested && !simultaneous_use_support) - || !mutable_support || BasicMutableCommandBufferTest::Skip(); + || !mutable_support; } - //-------------------------------------------------------------------------- cl_int Run() override { cl_int error = CL_SUCCESS; @@ -154,35 +165,32 @@ struct OutOfOrderTest : public BasicMutableCommandBufferTest return CL_SUCCESS; } - //-------------------------------------------------------------------------- cl_int RecordCommandBuffer() { cl_sync_point_khr sync_points[2]; const cl_int pattern = pattern_pri; - cl_int error = - clCommandFillBufferKHR(out_of_order_command_buffer, nullptr, in_mem, - &pattern, sizeof(cl_int), 0, data_size(), 0, - nullptr, &sync_points[0], nullptr); + cl_int error = clCommandFillBufferKHR( + work_command_buffer, nullptr, in_mem, &pattern, sizeof(cl_int), 0, + data_size(), 0, nullptr, &sync_points[0], nullptr); test_error(error, "clCommandFillBufferKHR failed"); - error = clCommandFillBufferKHR(out_of_order_command_buffer, nullptr, - out_mem, &overwritten_pattern, - sizeof(cl_int), 0, data_size(), 0, - nullptr, &sync_points[1], nullptr); + error = clCommandFillBufferKHR(work_command_buffer, nullptr, out_mem, + &overwritten_pattern, sizeof(cl_int), 0, + data_size(), 0, nullptr, &sync_points[1], + nullptr); test_error(error, "clCommandFillBufferKHR failed"); error = clCommandNDRangeKernelKHR( - out_of_order_command_buffer, nullptr, nullptr, kernel, 1, nullptr, + work_command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements, nullptr, 2, sync_points, nullptr, &command); test_error(error, "clCommandNDRangeKernelKHR failed"); - error = clFinalizeCommandBufferKHR(out_of_order_command_buffer); + error = clFinalizeCommandBufferKHR(work_command_buffer); test_error(error, "clFinalizeCommandBufferKHR failed"); return CL_SUCCESS; } - //-------------------------------------------------------------------------- cl_int RunSingle() { cl_int error; @@ -190,14 +198,14 @@ struct OutOfOrderTest : public BasicMutableCommandBufferTest error = RecordCommandBuffer(); test_error(error, "RecordCommandBuffer failed"); - error = clEnqueueCommandBufferKHR( - 0, nullptr, out_of_order_command_buffer, 0, nullptr, &single_event); + error = clEnqueueCommandBufferKHR(0, nullptr, work_command_buffer, 0, + nullptr, &single_event); test_error(error, "clEnqueueCommandBufferKHR failed"); std::vector output_data(num_elements); - error = clEnqueueReadBuffer(out_of_order_queue, out_mem, CL_TRUE, 0, - data_size(), output_data.data(), 1, - &single_event, nullptr); + error = + clEnqueueReadBuffer(work_queue, out_mem, CL_TRUE, 0, data_size(), + output_data.data(), 1, &single_event, nullptr); test_error(error, "clEnqueueReadBuffer failed"); for (size_t i = 0; i < num_elements; i++) @@ -235,15 +243,15 @@ struct OutOfOrderTest : public BasicMutableCommandBufferTest &dispatch_config }; - error = clUpdateMutableCommandsKHR(out_of_order_command_buffer, - &mutable_config); + error = + clUpdateMutableCommandsKHR(work_command_buffer, &mutable_config); test_error(error, "clUpdateMutableCommandsKHR failed"); - error = clEnqueueCommandBufferKHR( - 0, nullptr, out_of_order_command_buffer, 0, nullptr, &single_event); + error = clEnqueueCommandBufferKHR(0, nullptr, work_command_buffer, 0, + nullptr, &single_event); test_error(error, "clEnqueueCommandBufferKHR failed"); - error = clEnqueueReadBuffer(out_of_order_queue, new_out_mem, CL_TRUE, 0, + error = clEnqueueReadBuffer(work_queue, new_out_mem, CL_TRUE, 0, data_size(), output_data.data(), 1, &single_event, nullptr); test_error(error, "clEnqueueReadBuffer failed"); @@ -256,38 +264,35 @@ struct OutOfOrderTest : public BasicMutableCommandBufferTest return CL_SUCCESS; } - //-------------------------------------------------------------------------- cl_int RecordSimultaneousCommandBuffer() { cl_sync_point_khr sync_points[2]; // for both simultaneous passes this call will fill entire in_mem buffer cl_int error = clCommandFillBufferKHR( - out_of_order_command_buffer, nullptr, in_mem, &pattern_pri, - sizeof(cl_int), 0, data_size() * buffer_size_multiplier, 0, nullptr, + work_command_buffer, nullptr, in_mem, &pattern_pri, sizeof(cl_int), + 0, data_size() * buffer_size_multiplier, 0, nullptr, &sync_points[0], nullptr); test_error(error, "clCommandFillBufferKHR failed"); // to avoid overwriting the entire result buffer instead of filling // only relevant part this additional kernel was introduced - error = clCommandNDRangeKernelKHR(out_of_order_command_buffer, nullptr, - nullptr, kernel_fill, 1, nullptr, - &num_elements, nullptr, 0, nullptr, - &sync_points[1], &command); + error = clCommandNDRangeKernelKHR( + work_command_buffer, nullptr, nullptr, kernel_fill, 1, nullptr, + &num_elements, nullptr, 0, nullptr, &sync_points[1], &command); test_error(error, "clCommandNDRangeKernelKHR failed"); error = clCommandNDRangeKernelKHR( - out_of_order_command_buffer, nullptr, nullptr, kernel, 1, nullptr, + work_command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements, nullptr, 2, sync_points, nullptr, &command); test_error(error, "clCommandNDRangeKernelKHR failed"); - error = clFinalizeCommandBufferKHR(out_of_order_command_buffer); + error = clFinalizeCommandBufferKHR(work_command_buffer); test_error(error, "clFinalizeCommandBufferKHR failed"); return CL_SUCCESS; } - //-------------------------------------------------------------------------- struct SimulPassData { cl_int offset; @@ -296,7 +301,6 @@ struct OutOfOrderTest : public BasicMutableCommandBufferTest clEventWrapper wait_events[3]; }; - //-------------------------------------------------------------------------- cl_int EnqueueSimultaneousPass(SimulPassData& pd) { cl_int error = CL_SUCCESS; @@ -310,19 +314,19 @@ struct OutOfOrderTest : public BasicMutableCommandBufferTest // filling offset buffer must wait for previous pass completeness error = clEnqueueFillBuffer( - out_of_order_queue, off_mem, &pd.offset, sizeof(cl_int), 0, - sizeof(cl_int), (wait_pass_event != nullptr ? 1 : 0), + work_queue, off_mem, &pd.offset, sizeof(cl_int), 0, sizeof(cl_int), + (wait_pass_event != nullptr ? 1 : 0), (wait_pass_event != nullptr ? &wait_pass_event : nullptr), &pd.wait_events[1]); test_error(error, "clEnqueueFillBuffer failed"); // command buffer execution must wait for two wait-events - error = clEnqueueCommandBufferKHR( - 0, nullptr, out_of_order_command_buffer, 2, &pd.wait_events[0], - &pd.wait_events[2]); + error = + clEnqueueCommandBufferKHR(0, nullptr, work_command_buffer, 2, + &pd.wait_events[0], &pd.wait_events[2]); test_error(error, "clEnqueueCommandBufferKHR failed"); - error = clEnqueueReadBuffer(out_of_order_queue, out_mem, CL_FALSE, + error = clEnqueueReadBuffer(work_queue, out_mem, CL_FALSE, pd.offset * sizeof(cl_int), data_size(), pd.output_buffer.data(), 1, &pd.wait_events[2], nullptr); @@ -358,17 +362,17 @@ struct OutOfOrderTest : public BasicMutableCommandBufferTest &dispatch_config }; - error = clUpdateMutableCommandsKHR(out_of_order_command_buffer, - &mutable_config); + error = + clUpdateMutableCommandsKHR(work_command_buffer, &mutable_config); test_error(error, "clUpdateMutableCommandsKHR failed"); // command buffer execution must wait for two wait-events - error = clEnqueueCommandBufferKHR( - 0, nullptr, out_of_order_command_buffer, 2, &pd.wait_events[0], - &pd.wait_events[2]); + error = + clEnqueueCommandBufferKHR(0, nullptr, work_command_buffer, 2, + &pd.wait_events[0], &pd.wait_events[2]); test_error(error, "clEnqueueCommandBufferKHR failed"); - error = clEnqueueReadBuffer(out_of_order_queue, new_out_mem, CL_FALSE, + error = clEnqueueReadBuffer(work_queue, new_out_mem, CL_FALSE, pd.offset * sizeof(cl_int), data_size(), pd.output_buffer.data(), 1, &pd.wait_events[2], nullptr); @@ -377,7 +381,6 @@ struct OutOfOrderTest : public BasicMutableCommandBufferTest return CL_SUCCESS; } - //-------------------------------------------------------------------------- cl_int RunSimultaneous() { cl_int error = RecordSimultaneousCommandBuffer(); @@ -401,7 +404,7 @@ struct OutOfOrderTest : public BasicMutableCommandBufferTest error = clSetUserEventStatus(user_event, CL_COMPLETE); test_error(error, "clSetUserEventStatus failed"); - error = clFinish(out_of_order_queue); + error = clFinish(work_queue); test_error(error, "clFinish failed"); // verify the result buffers @@ -417,9 +420,8 @@ struct OutOfOrderTest : public BasicMutableCommandBufferTest return CL_SUCCESS; } - //-------------------------------------------------------------------------- - clCommandQueueWrapper out_of_order_queue; - clCommandBufferWrapper out_of_order_command_buffer; + clCommandQueueWrapper work_queue; + clCommandBufferWrapper work_command_buffer; clEventWrapper user_event; clEventWrapper single_event; @@ -429,10 +431,177 @@ struct OutOfOrderTest : public BasicMutableCommandBufferTest clProgramWrapper program_fill; const size_t test_global_work_size = 3 * sizeof(cl_int); - cl_mutable_command_khr command = nullptr; + const cl_int pattern_pri = 42; const cl_int overwritten_pattern = 0xACDC; + cl_mutable_command_khr command; +}; + +struct CrossQueueSimultaneousMutableDispatchTest + : public BasicMutableCommandBufferTest +{ + CrossQueueSimultaneousMutableDispatchTest(cl_device_id device, + cl_context context, + cl_command_queue queue) + : BasicMutableCommandBufferTest(device, context, queue), + queue_sec(nullptr), command(nullptr) + { + simultaneous_use_requested = true; + } + + cl_int SetUpKernel() override + { + const char* kernel_str = + R"( + __kernel void fill(int pattern, __global int* out) + { + size_t id = get_global_id(0); + out[id] = pattern; + })"; + + cl_int error = create_single_kernel_helper_create_program( + context, &program, 1, &kernel_str); + test_error(error, "Failed to create program with source"); + + error = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr); + test_error(error, "Failed to build program"); + + kernel = clCreateKernel(program, "fill", &error); + test_error(error, "Failed to create copy kernel"); + + return CL_SUCCESS; + } + + cl_int SetUpKernelArgs() override + { + cl_int error = CL_SUCCESS; + out_mem = clCreateBuffer(context, CL_MEM_WRITE_ONLY, data_size(), + nullptr, &error); + test_error(error, "clCreateBuffer failed"); + + error = clSetKernelArg(kernel, 0, sizeof(cl_int), &pattern_pri); + test_error(error, "clSetKernelArg failed"); + + error = clSetKernelArg(kernel, 1, sizeof(out_mem), &out_mem); + test_error(error, "clSetKernelArg failed"); + + return CL_SUCCESS; + } + + cl_int SetUp(int elements) override + { + cl_int error = BasicMutableCommandBufferTest::SetUp(elements); + test_error(error, "BasicMutableCommandBufferTest::SetUp failed"); + + queue_sec = clCreateCommandQueue(context, device, 0, &error); + test_error(error, "Unable to create command queue to test with"); + + return CL_SUCCESS; + } + + bool Skip() override + { + if (BasicMutableCommandBufferTest::Skip()) return true; + + cl_mutable_dispatch_fields_khr mutable_capabilities = { 0 }; + + bool mutable_support = + !clGetDeviceInfo( + device, CL_DEVICE_MUTABLE_DISPATCH_CAPABILITIES_KHR, + sizeof(mutable_capabilities), &mutable_capabilities, nullptr) + && mutable_capabilities & CL_MUTABLE_DISPATCH_ARGUMENTS_KHR; + + return !simultaneous_use_support || !mutable_support; + } + + cl_int Run() override + { + // record command buffer + cl_int pattern = 0; + cl_int error = clCommandFillBufferKHR( + command_buffer, nullptr, out_mem, &pattern, sizeof(cl_int), 0, + data_size(), 0, nullptr, nullptr, nullptr); + test_error(error, "clCommandFillBufferKHR failed"); + + cl_ndrange_kernel_command_properties_khr props[] = { + CL_MUTABLE_DISPATCH_UPDATABLE_FIELDS_KHR, + CL_MUTABLE_DISPATCH_ARGUMENTS_KHR, 0 + }; + + error = clCommandNDRangeKernelKHR( + command_buffer, nullptr, props, kernel, 1, nullptr, &num_elements, + nullptr, 0, nullptr, nullptr, &command); + test_error(error, "clCommandNDRangeKernelKHR failed"); + + error = clFinalizeCommandBufferKHR(command_buffer); + test_error(error, "clFinalizeCommandBufferKHR failed"); + + // enqueue command buffer to default queue + error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0, + nullptr, nullptr); + test_error(error, "clEnqueueCommandBufferKHR failed"); + + // update mutable parameters + clMemWrapper new_out_mem = clCreateBuffer(context, CL_MEM_WRITE_ONLY, + data_size(), nullptr, &error); + test_error(error, "clCreateBuffer failed"); + + cl_mutable_dispatch_arg_khr arg_0{ 0, sizeof(cl_int), &pattern_sec }; + cl_mutable_dispatch_arg_khr arg_1{ 1, sizeof(new_out_mem), + &new_out_mem }; + cl_mutable_dispatch_arg_khr args[] = { arg_0, arg_1 }; + + cl_mutable_dispatch_config_khr dispatch_config{ + CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR, + nullptr, + command, + 2 /* num_args */, + 0 /* num_svm_arg */, + 0 /* num_exec_infos */, + 0 /* work_dim - 0 means no change to dimensions */, + args /* 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_mutable_base_config_khr mutable_config{ + CL_STRUCTURE_TYPE_MUTABLE_BASE_CONFIG_KHR, nullptr, 1, + &dispatch_config + }; + + error = clUpdateMutableCommandsKHR(command_buffer, &mutable_config); + test_error(error, "clUpdateMutableCommandsKHR failed"); + + // enqueue command buffer to non-default queue + error = clEnqueueCommandBufferKHR(1, &queue_sec, command_buffer, 0, + nullptr, nullptr); + test_error(error, "clEnqueueCommandBufferKHR failed"); + + error = clFinish(queue_sec); + test_error(error, "clFinish failed"); + + // read result of command buffer execution + std::vector output_data(num_elements); + error = + clEnqueueReadBuffer(queue_sec, new_out_mem, CL_TRUE, 0, data_size(), + output_data.data(), 0, nullptr, nullptr); + test_error(error, "clEnqueueReadBuffer failed"); + + // verify the result + for (size_t i = 0; i < num_elements; i++) + { + CHECK_VERIFICATION_ERROR(pattern_sec, output_data[i], i); + } + + return CL_SUCCESS; + } + + clCommandQueueWrapper queue_sec; const cl_int pattern_pri = 42; + const cl_int pattern_sec = 0xACDC; + cl_mutable_command_khr command; }; } // anonymous namespace @@ -440,8 +609,8 @@ struct OutOfOrderTest : public BasicMutableCommandBufferTest int test_mutable_dispatch_out_of_order(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) { - return MakeAndRunTest>(device, context, queue, - num_elements); + return MakeAndRunTest>( + device, context, queue, num_elements); } int test_mutable_dispatch_simultaneous_out_of_order(cl_device_id device, @@ -449,6 +618,24 @@ int test_mutable_dispatch_simultaneous_out_of_order(cl_device_id device, cl_command_queue queue, int num_elements) { - return MakeAndRunTest>(device, context, queue, - num_elements); + return MakeAndRunTest>( + device, context, queue, num_elements); +} + +int test_mutable_dispatch_simultaneous_in_order(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements) +{ + return MakeAndRunTest>( + device, context, queue, num_elements); +} + +int test_mutable_dispatch_simultaneous_cross_queue(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements) +{ + return MakeAndRunTest( + device, context, queue, num_elements); } diff --git a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/procs.h b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/procs.h index 5991f24a8e..ca5ab1ffe3 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/procs.h +++ b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/procs.h @@ -94,6 +94,13 @@ extern int test_mutable_dispatch_out_of_order(cl_device_id device, extern int test_mutable_dispatch_simultaneous_out_of_order( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements); +extern int test_mutable_dispatch_simultaneous_in_order(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_mutable_dispatch_simultaneous_cross_queue( + cl_device_id device, cl_context context, cl_command_queue queue, + int num_elements); extern int test_mutable_dispatch_global_size(cl_device_id device, cl_context context, cl_command_queue queue, From da4a30d394c04667178d74902092d2120963c863 Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Tue, 19 Mar 2024 16:42:42 +0100 Subject: [PATCH 3/3] Added new test to cover multiple commands dispatch in one enqueued update (#1919) * Added new test to cover multiple comands dispatch in one enqueued update According to issue description #1481 * clang format correction * Few minor corrections * cosmetic corrections --- .../CMakeLists.txt | 1 + .../main.cpp | 1 + .../mutable_command_multiple_dispatches.cpp | 217 ++++++++++++++++++ .../procs.h | 4 + 4 files changed, 223 insertions(+) create mode 100644 test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_multiple_dispatches.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 16f847d0e7..3b3af6980a 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 @@ -10,6 +10,7 @@ set(${MODULE_NAME}_SOURCES mutable_command_local_size.cpp mutable_command_global_offset.cpp mutable_command_full_dispatch.cpp + mutable_command_multiple_dispatches.cpp ../basic_command_buffer.cpp ) diff --git a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/main.cpp b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/main.cpp index e6af2898c9..9a09685e42 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/main.cpp +++ b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/main.cpp @@ -27,6 +27,7 @@ test_definition test_list[] = { ADD_TEST(mutable_command_info_local_work_size), ADD_TEST(mutable_command_info_global_work_size), ADD_TEST(mutable_command_full_dispatch), + ADD_TEST(mutable_command_multiple_dispatches), ADD_TEST(mutable_dispatch_image_1d_arguments), ADD_TEST(mutable_dispatch_image_2d_arguments), ADD_TEST(mutable_dispatch_out_of_order), diff --git a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_multiple_dispatches.cpp b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_multiple_dispatches.cpp new file mode 100644 index 0000000000..0911472003 --- /dev/null +++ b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_multiple_dispatches.cpp @@ -0,0 +1,217 @@ +// +// Copyright (c) 2024 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 +#include "mutable_command_basic.h" + +#include +#include + +#include + +namespace { + +//////////////////////////////////////////////////////////////////////////////// +// command buffer with multiple command handles dispatch test + +struct MultipleCommandsDispatch : BasicMutableCommandBufferTest +{ + MultipleCommandsDispatch(cl_device_id device, cl_context context, + cl_command_queue queue) + : BasicMutableCommandBufferTest(device, context, queue), + command_pri(nullptr), command_sec(nullptr) + { + simultaneous_use_requested = false; + } + + 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 capabillity + return !mutable_support; + } + + // setup default and fill kernels program + cl_int SetUpKernel() override + { + // default command buffer kernel + cl_int error = BasicCommandBufferTest::SetUpKernel(); + test_error(error, "BasicCommandBufferTest::SetUpKernel failed"); + + // fill command buffer kernel + const char *kernel_fill_str = + R"( + __kernel void fill(int pattern, __global int *dst) + { + size_t gid = get_global_id(0); + dst[gid] = pattern; + })"; + + error = create_single_kernel_helper_create_program( + context, &program_fill, 1, &kernel_fill_str); + test_error(error, "Failed to create program with source"); + + error = + clBuildProgram(program_fill, 1, &device, nullptr, nullptr, nullptr); + test_error(error, "Failed to build program"); + + kernel_fill = clCreateKernel(program_fill, "fill", &error); + test_error(error, "Failed to create copy kernel"); + + return CL_SUCCESS; + } + + // setup kernel arguments for both default and fill kernels + cl_int SetUpKernelArgs() override + { + // arguments for default kernel + cl_int error = BasicCommandBufferTest::SetUpKernelArgs(); + test_error(error, "BasicCommandBufferTest::SetUpKernelArgs failed"); + + // fill kernel applies pattern for input data of default kernel + error = clSetKernelArg(kernel_fill, 0, sizeof(cl_int), &pattern_pri); + test_error(error, "clSetKernelArg failed"); + + error = clSetKernelArg(kernel_fill, 1, sizeof(in_mem), &in_mem); + test_error(error, "clSetKernelArg failed"); + + return CL_SUCCESS; + } + + // Check the results of command buffer execution + bool verify_result(const cl_mem &buffer, const cl_int pattern) + { + cl_int error = CL_SUCCESS; + std::vector data(num_elements); + 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] != pattern) + { + log_error("Modified verification failed at index %zu: Got %d, " + "wanted %d\n", + i, data[i], pattern); + return false; + } + } + + return true; + } + + // run command buffer with multiple command dispatches test + cl_int Run() override + { + // record fill kernel and collect first mutable command handle + cl_int error = clCommandNDRangeKernelKHR( + command_buffer, nullptr, nullptr, kernel_fill, 1, nullptr, + &num_elements, nullptr, 0, nullptr, nullptr, &command_pri); + test_error(error, "clCommandNDRangeKernelKHR failed"); + + // record default kernel and collect second mutable command handle + error = clCommandNDRangeKernelKHR( + command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements, + nullptr, 0, nullptr, nullptr, &command_sec); + test_error(error, "clCommandNDRangeKernelKHR failed"); + + error = clFinalizeCommandBufferKHR(command_buffer); + test_error(error, "clFinalizeCommandBufferKHR failed"); + + error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0, + nullptr, nullptr); + test_error(error, "clEnqueueCommandBufferKHR failed"); + + error = clFinish(queue); + test_error(error, "clFinish failed"); + + // check the results of the initial execution + if (!verify_result(out_mem, pattern_pri)) return TEST_FAIL; + + // new output buffer for default kernel + clMemWrapper new_out_mem = clCreateBuffer(context, CL_MEM_WRITE_ONLY, + data_size(), nullptr, &error); + test_error(error, "clCreateBuffer failed"); + + // apply dispatch for mutable arguments of both fill and default kernels + cl_mutable_dispatch_arg_khr arg_pri{ 0, sizeof(cl_int), &pattern_sec }; + cl_mutable_dispatch_arg_khr args_pri[] = { arg_pri }; + + cl_mutable_dispatch_arg_khr arg_sec{ 1, sizeof(new_out_mem), + &new_out_mem }; + cl_mutable_dispatch_arg_khr args_sec[] = { arg_sec }; + + // modify two mutable parameters, each one with separate handle + cl_mutable_dispatch_config_khr dispatch_config[] = { + { CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR, nullptr, + command_pri, 1, 0, 0, 0, args_pri, nullptr, nullptr, nullptr, + nullptr, nullptr }, + { CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR, nullptr, + command_sec, 1, 0, 0, 0, args_sec, nullptr, nullptr, nullptr, + nullptr, nullptr }, + }; + + cl_mutable_base_config_khr mutable_config{ + CL_STRUCTURE_TYPE_MUTABLE_BASE_CONFIG_KHR, nullptr, 2, + dispatch_config + }; + + error = clUpdateMutableCommandsKHR(command_buffer, &mutable_config); + test_error(error, "clUpdateMutableCommandsKHR failed"); + + // repeat execution of modified command buffer + error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0, + nullptr, nullptr); + test_error(error, "clEnqueueCommandBufferKHR failed"); + + error = clFinish(queue); + test_error(error, "clFinish failed"); + + // check the results of the modified execution + if (!verify_result(new_out_mem, pattern_sec)) return TEST_FAIL; + + return TEST_PASS; + } + + // mutable dispatch test attributes + cl_mutable_command_khr command_pri; + cl_mutable_command_khr command_sec; + + clKernelWrapper kernel_fill; + clProgramWrapper program_fill; + + const cl_int pattern_pri = 0xACDC; + const cl_int pattern_sec = 0xDEAD; +}; + +} + +int test_mutable_command_multiple_dispatches(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements) +{ + return MakeAndRunTest(device, context, queue, + num_elements); +} diff --git a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/procs.h b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/procs.h index ca5ab1ffe3..bce6ad8311 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/procs.h +++ b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/procs.h @@ -117,5 +117,9 @@ extern int test_mutable_command_full_dispatch(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements); +extern int test_mutable_command_multiple_dispatches(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements); #endif /*_CL_KHR_COMMAND_BUFFER_MUTABLE_DISPATCH_PROCS_H*/