Add local size tests for cl_khr_command_buffer_mutable_dispatch. (#1745)

* Add local size tests for cl_khr_command_buffer_mutable_dispatch.

Signed-off-by: Paweł Jastrzębski <p.k.jastrzebski@gmail.com>

* Add kernel with observable output.

We should check that's some observable output from the kernel
as a result of the change to local work size, not just that
clGetMutableCommandInfoKHR has been updated. For example,
getting every work-item to call get_local_size() inside of
the kernel and writing it to a buffer, then reading the buffer
after the command-buffer enqueue has finished and check it
matches what we expect.

Signed-off-by: Paweł Jastrzębski <p.k.jastrzebski@gmail.com>

* Fix review comments.

    Applied review comments for mutable dispatch local size test:
    - clFinish to ensure command-buffer has finished executing for calling clUpdateMutableCommandsKHR
    - Change variable and constant names for local size

Applied review comments for mutable dispatch global arguments test:
- clFinish to ensure command-buffer has finished executing for calling clUpdateMutableCommandsKHR

Signed-off-by: Paweł Jastrzębski <p.k.jastrzebski@gmail.com>

* Fix review comments.

Changes made:
- Fix skip conditions
- Remove obsolete variable
- Replace a variable with a constant

Signed-off-by: Paweł Jastrzębski <p.k.jastrzebski@gmail.com>

* Remove explicit base class call.

Signed-off-by: Paweł Jastrzębski <p.k.jastrzebski@gmail.com>

* Fix constant magic number.

Signed-off-by: Paweł Jastrzębski <p.k.jastrzebski@gmail.com>

* Update global size and local size to meet the spec requirements.

Make sure work-groups number is not increased after update of command-buffer.

Signed-off-by: Paweł Jastrzębski <p.k.jastrzebski@gmail.com>

* Remove uneeded includes.

Signed-off-by: Paweł Jastrzębski <p.k.jastrzebski@gmail.com>

---------

Signed-off-by: Paweł Jastrzębski <p.k.jastrzebski@gmail.com>

test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/CMakeLists.txt

@@ -3,6 +3,7 @@ set(MODULE_NAME CL_KHR_MUTABLE_DISPATCH)
 set(${MODULE_NAME}_SOURCES
     main.cpp
     mutable_command_info.cpp
+    mutable_command_local_size.cpp
     mutable_command_global_offset.cpp
     ../basic_command_buffer.cpp
 )

test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/main.cpp

@@ -26,6 +26,7 @@ test_definition test_list[] = {
     ADD_TEST(mutable_command_info_global_work_offset),
     ADD_TEST(mutable_command_info_local_work_size),
     ADD_TEST(mutable_command_info_global_work_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_local_size.cpp

@@ -0,0 +1,174 @@
+//
+// Copyright (c) 2022 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 "typeWrappers.h"
+#include "procs.h"
+#include "testHarness.h"
+#include "mutable_command_basic.h"
+#include <vector>
+
+#include <CL/cl.h>
+#include <CL/cl_ext.h>
+
+////////////////////////////////////////////////////////////////////////////////
+// mutable dispatch tests which handle following cases:
+//
+// CL_MUTABLE_DISPATCH_LOCAL_WORK_SIZE_KHR
+
+struct MutableDispatchLocalSize : public InfoMutableCommandBufferTest
+{
+    using InfoMutableCommandBufferTest::InfoMutableCommandBufferTest;
+
+    MutableDispatchLocalSize(cl_device_id device, cl_context context,
+                             cl_command_queue queue)
+        : InfoMutableCommandBufferTest(device, context, queue)
+    {}
+
+    bool Skip() override
+    {
+        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_LOCAL_SIZE_KHR;
+
+        return !mutable_support || InfoMutableCommandBufferTest::Skip();
+    }
+
+    cl_int Run() override
+    {
+        const char *local_size_kernel =
+            R"(
+                __kernel void sample_test(__global int *dst)
+            {
+                size_t tid = get_global_id(0);
+                dst[tid] = get_local_size(0);
+            })";
+
+        cl_int error = create_single_kernel_helper(
+            context, &program, &kernel, 1, &local_size_kernel, "sample_test");
+        test_error(error, "Creating kernel failed");
+
+        clMemWrapper stream;
+        stream = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeToAllocate,
+                                nullptr, &error);
+        test_error(error, "Creating test array failed");
+
+        /* Set the arguments */
+        error = clSetKernelArg(kernel, 0, sizeof(cl_mem), &stream);
+        test_error(error, "Unable to set indexed kernel arguments");
+
+        error = clCommandNDRangeKernelKHR(
+            command_buffer, nullptr, nullptr, kernel, 1, nullptr,
+            &global_work_size, &local_work_size, 0, nullptr, nullptr, &command);
+        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.");
+
+        cl_mutable_dispatch_config_khr dispatch_config{
+            CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR,
+            nullptr,
+            command,
+            0 /* num_args */,
+            0 /* num_svm_arg */,
+            0 /* num_exec_infos */,
+            0 /* work_dim - 0 means no change to dimensions */,
+            nullptr /* arg_list */,
+            nullptr /* arg_svm_list - nullptr means no change*/,
+            nullptr /* exec_info_list */,
+            nullptr /* global_work_offset */,
+            &update_global_size /* global_work_size */,
+            &update_local_size /* 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");
+
+        error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
+                                          nullptr, nullptr);
+        test_error(error, "clEnqueueCommandBufferKHR failed");
+
+        error = clGetMutableCommandInfoKHR(
+            command, CL_MUTABLE_DISPATCH_LOCAL_WORK_SIZE_KHR,
+            sizeof(info_local_size), &info_local_size, nullptr);
+        test_error(error, "clGetMutableCommandInfoKHR failed");
+
+        if (info_local_size != update_local_size)
+        {
+            log_error("ERROR: Wrong size returned from "
+                      "clGetMutableCommandInfoKHR.");
+            return TEST_FAIL;
+        }
+
+        std::vector<cl_int> resultData;
+        resultData.resize(num_elements);
+
+        error = clEnqueueReadBuffer(queue, stream, CL_TRUE, 0, sizeToAllocate,
+                                    resultData.data(), 0, nullptr, nullptr);
+        test_error(error, "clEnqueueReadBuffer failed");
+
+        for (size_t i = 0; i < num_elements; i++)
+            if (i < update_global_size && update_local_size != resultData[i])
+            {
+                log_error("Data failed to verify: update_local_size != "
+                          "resultData[%d]=%d\n",
+                          i, resultData[i]);
+                return TEST_FAIL;
+            }
+            else if (i >= update_global_size
+                     && local_work_size != resultData[i])
+            {
+                log_error("Data failed to verify: update_local_size != "
+                          "resultData[%d]=%d\n",
+                          i, resultData[i]);
+                return TEST_FAIL;
+            }
+
+        return CL_SUCCESS;
+    }
+
+    size_t info_local_size = 0;
+    const size_t global_work_size = 16;
+    const size_t local_work_size = 8;
+    const size_t update_global_size = 8;
+    const size_t update_local_size = 4;
+    const size_t sizeToAllocate = 64;
+    const size_t num_elements = sizeToAllocate / sizeof(cl_int);
+
+    cl_mutable_command_khr command = nullptr;
+};
+
+int test_mutable_dispatch_local_size(cl_device_id device, cl_context context,
+                                     cl_command_queue queue, int num_elements)
+{
+    return MakeAndRunTest<MutableDispatchLocalSize>(device, context, queue,
+                                                    num_elements);
+}

test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/procs.h

@@ -59,6 +59,10 @@ extern int test_mutable_command_info_global_work_size(cl_device_id device,
                                                       cl_context context,
                                                       cl_command_queue queue,
                                                       int num_elements);
+extern int test_mutable_dispatch_local_size(cl_device_id device,
+                                            cl_context context,
+                                            cl_command_queue queue,
+                                            int num_elements);
 extern int test_mutable_dispatch_global_offset(cl_device_id device,
                                                cl_context context,
                                                cl_command_queue queue,