Merge branch 'main' into mutable_command_iterative_arg_update

test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/CMakeLists.txt
test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/main.cpp
test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/procs.h

.github/dependabot.yml

@@ -0,0 +1,10 @@
+version: 2
+updates:
+  - package-ecosystem: "github-actions"
+    directory: "/"
+    schedule:
+      interval: "monthly"
+    groups:
+      github-actions:
+        patterns:
+          - "*"

test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/CMakeLists.txt

@@ -5,12 +5,13 @@ 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
     mutable_command_full_dispatch.cpp
-    mutable_command_iterative_arg_update.cpp
+    mutable_command_multiple_dispatches.cpp
+	mutable_command_iterative_arg_update.cpp
     ../basic_command_buffer.cpp
 )
 

test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/main.cpp

@@ -27,11 +27,14 @@ 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_command_iterative_arg_update),
     ADD_TEST(mutable_dispatch_image_1d_arguments),
     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),

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,
         };
 

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 <extensionHelpers.h>
+#include "mutable_command_basic.h"
+
+#include <CL/cl.h>
+#include <CL/cl_ext.h>
+
+#include <vector>
+
+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<cl_int> 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<MultipleCommandsDispatch>(device, context, queue,
+                                                    num_elements);
+}