Add global offset tests for cl_khr_command_buffer_mutable_dispatch. (#1743)

* Add global offset 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 there's some observable output from the kernel as a result of the change to global work offset, not just that clGetMutableCommandInfoKHR has been updated. E.g we could call get_global_offset() inside of the kernel, write something to a buffer based on that, and read the buffer after the command-buffer enqueue has finished. Signed-off-by: Paweł Jastrzębski <p.k.jastrzebski@gmail.com> * Fix review comments. Applied review comments for mutable dispatch global offset test: - clFinish to ensure command-buffer has finished executing for calling clUpdateMutableCommandsKHR - Change variable and constant names for global offset - Remove redundant return CL_SUCCESS 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> * Fix review comments. Changes made: - Remove explicit base class call - Fix constant magic number 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> * Fix clang-format. Signed-off-by: Paweł Jastrzębski <p.k.jastrzebski@gmail.com> * Fix condition for result check. Signed-off-by: Paweł Jastrzębski <p.k.jastrzebski@gmail.com> --------- Signed-off-by: Paweł Jastrzębski <p.k.jastrzebski@gmail.com>
2026-03-19 06:09:01 +00:00 · 2023-06-27 17:54:14 +02:00
parent fee6d6bb66
commit 56974a5858
6 changed files with 267 additions and 43 deletions
--- 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
@@ -3,6 +3,7 @@ set(MODULE_NAME CL_KHR_MUTABLE_DISPATCH)
 set(${MODULE_NAME}_SOURCES
    main.cpp
    mutable_command_info.cpp
+    mutable_command_global_offset.cpp
    ../basic_command_buffer.cpp
 )

--- 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
@@ -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_global_offset),
 };

 int main(int argc, const char *argv[])
--- 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
@@ -19,6 +19,17 @@
 #include "../basic_command_buffer.h"
 #include "../command_buffer_test_base.h"

+// If it is supported get the addresses of all the APIs here.
+#define GET_EXTENSION_ADDRESS(FUNC)                                            \
+    FUNC = reinterpret_cast<FUNC##_fn>(                                        \
+        clGetExtensionFunctionAddressForPlatform(platform, #FUNC));            \
+    if (FUNC == nullptr)                                                       \
+    {                                                                          \
+        log_error("ERROR: clGetExtensionFunctionAddressForPlatform failed"     \
+                  " with " #FUNC "\n");                                        \
+        return TEST_FAIL;                                                      \
+    }
+
 struct BasicMutableCommandBufferTest : BasicCommandBufferTest
 {
    BasicMutableCommandBufferTest(cl_device_id device, cl_context context,
@@ -84,24 +95,52 @@ struct BasicMutableCommandBufferTest : BasicCommandBufferTest
                            &platform, nullptr);
        test_error(error, "clGetDeviceInfo for CL_DEVICE_PLATFORM failed");

-        // If it is supported get the addresses of all the APIs here.
-#define GET_EXTENSION_ADDRESS(FUNC)                                            \
-    FUNC = reinterpret_cast<FUNC##_fn>(                                        \
-        clGetExtensionFunctionAddressForPlatform(platform, #FUNC));            \
-    if (FUNC == nullptr)                                                       \
-    {                                                                          \
-        log_error("ERROR: clGetExtensionFunctionAddressForPlatform failed"     \
-                  " with " #FUNC "\n");                                        \
-        return TEST_FAIL;                                                      \
+        GET_EXTENSION_ADDRESS(clUpdateMutableCommandsKHR);
+
+        return CL_SUCCESS;
    }
+
+    clUpdateMutableCommandsKHR_fn clUpdateMutableCommandsKHR = nullptr;
+
+    const char* kernelString = "__kernel void empty() {}";
+    const size_t global_work_size = 4 * 16;
+};
+
+struct InfoMutableCommandBufferTest : BasicMutableCommandBufferTest
+{
+    InfoMutableCommandBufferTest(cl_device_id device, cl_context context,
+                                 cl_command_queue queue)
+        : BasicMutableCommandBufferTest(device, context, queue)
+    {}
+
+    virtual cl_int SetUp(int elements) override
+    {
+        BasicMutableCommandBufferTest::SetUp(elements);
+
+        cl_int error = init_extension_functions();
+        test_error(error, "Unable to initialise extension functions");
+
+        return CL_SUCCESS;
+    }
+
+    cl_int init_extension_functions()
+    {
+        BasicCommandBufferTest::init_extension_functions();
+
+        cl_platform_id platform;
+        cl_int error =
+            clGetDeviceInfo(device, CL_DEVICE_PLATFORM, sizeof(cl_platform_id),
+                            &platform, nullptr);
+        test_error(error, "clGetDeviceInfo for CL_DEVICE_PLATFORM failed");
+
        GET_EXTENSION_ADDRESS(clGetMutableCommandInfoKHR);

        return CL_SUCCESS;
    }

    clGetMutableCommandInfoKHR_fn clGetMutableCommandInfoKHR = nullptr;
-    const char* kernelString = "__kernel void empty() {}";
-    const size_t global_work_size = 4 * sizeof(cl_int);
 };

-#endif // CL_KHR_MUTABLE_COMMAND_BASIC_H
+#undef GET_EXTENSION_ADDRESS
+
+#endif //_CL_KHR_MUTABLE_COMMAND_BASIC_H
--- a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_global_offset.cpp
+++ b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_global_offset.cpp
@@ -0,0 +1,179 @@
+//
+// 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 "imageHelpers.h"
+#include <vector>
+#include <iostream>
+#include <random>
+#include <cstring>
+#include <algorithm>
+#include <memory>
+#include "mutable_command_basic.h"
+
+#include <CL/cl.h>
+#include <CL/cl_ext.h>
+
+////////////////////////////////////////////////////////////////////////////////
+// mutable dispatch tests which handle following cases:
+//
+// CL_MUTABLE_DISPATCH_GLOBAL_WORK_OFFSET_KHR
+
+struct MutableDispatchGlobalOffset : InfoMutableCommandBufferTest
+{
+    using InfoMutableCommandBufferTest::InfoMutableCommandBufferTest;
+
+    MutableDispatchGlobalOffset(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_GLOBAL_OFFSET_KHR;
+
+        return !mutable_support || InfoMutableCommandBufferTest::Skip();
+    }
+
+    cl_int Run() override
+    {
+        const char *global_offset_kernel =
+            R"(
+                __kernel void sample_test(__global int *dst)
+            {
+                size_t tid = get_global_id(0);
+                dst[tid] = get_global_offset(0);
+            })";
+
+        cl_int error =
+            create_single_kernel_helper(context, &program, &kernel, 1,
+                                        &global_offset_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, nullptr, 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 */,
+            &update_global_offset /* 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");
+
+        error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
+                                          nullptr, nullptr);
+        test_error(error, "clEnqueueCommandBufferKHR failed");
+
+        error = clGetMutableCommandInfoKHR(
+            command, CL_MUTABLE_DISPATCH_GLOBAL_WORK_OFFSET_KHR,
+            sizeof(info_global_offset), &info_global_offset, nullptr);
+        test_error(error, "clGetMutableCommandInfoKHR failed");
+
+        if (info_global_offset != update_global_offset)
+        {
+            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_offset && 0 != resultData[i])
+            {
+                log_error("Data failed to verify: update_global_offset != "
+                          "resultData[%d]=%d\n",
+                          i, resultData[i]);
+                return TEST_FAIL;
+            }
+            else if (i >= update_global_offset
+                     && update_global_offset != resultData[i])
+            {
+                log_error("Data failed to verify: update_global_offset != "
+                          "resultData[%d]=%d\n",
+                          i, resultData[i]);
+                return TEST_FAIL;
+            }
+        return CL_SUCCESS;
+    }
+
+    size_t info_global_offset = 0;
+    const size_t update_global_offset = 3;
+    const size_t sizeToAllocate =
+        (global_work_size + update_global_offset) * sizeof(cl_int);
+    const size_t num_elements = sizeToAllocate / sizeof(cl_int);
+    cl_mutable_command_khr command = nullptr;
+};
+
+int test_mutable_dispatch_global_offset(cl_device_id device, cl_context context,
+                                        cl_command_queue queue,
+                                        int num_elements)
+{
+
+    return MakeAndRunTest<MutableDispatchGlobalOffset>(device, context, queue,
+                                                       num_elements);
+}
--- a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_info.cpp
+++ b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_info.cpp
@@ -42,13 +42,13 @@
 // CL_MUTABLE_DISPATCH_LOCAL_WORK_SIZE_KHR
 // CL_MUTABLE_COMMAND_COMMAND_TYPE_KHR

-struct InfoDeviceQuery : public BasicMutableCommandBufferTest
+struct InfoDeviceQuery : public InfoMutableCommandBufferTest
 {
-    using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest;
+    using InfoMutableCommandBufferTest::InfoMutableCommandBufferTest;

    InfoDeviceQuery(cl_device_id device, cl_context context,
                    cl_command_queue queue)
-        : BasicMutableCommandBufferTest(device, context, queue)
+        : InfoMutableCommandBufferTest(device, context, queue)
    {}

    cl_int Run() override
@@ -71,12 +71,12 @@ struct InfoDeviceQuery : public BasicMutableCommandBufferTest
    }
 };

-struct InfoBuffer : public BasicMutableCommandBufferTest
+struct InfoBuffer : public InfoMutableCommandBufferTest
 {
-    using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest;
+    using InfoMutableCommandBufferTest::InfoMutableCommandBufferTest;

    InfoBuffer(cl_device_id device, cl_context context, cl_command_queue queue)
-        : BasicMutableCommandBufferTest(device, context, queue)
+        : InfoMutableCommandBufferTest(device, context, queue)
    {}

    cl_int Run() override
@@ -108,13 +108,13 @@ struct InfoBuffer : public BasicMutableCommandBufferTest
    cl_mutable_command_khr command = nullptr;
 };

-struct PropertiesArray : public BasicMutableCommandBufferTest
+struct PropertiesArray : public InfoMutableCommandBufferTest
 {
-    using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest;
+    using InfoMutableCommandBufferTest::InfoMutableCommandBufferTest;

    PropertiesArray(cl_device_id device, cl_context context,
                    cl_command_queue queue)
-        : BasicMutableCommandBufferTest(device, context, queue)
+        : InfoMutableCommandBufferTest(device, context, queue)
    {}

    cl_int Run() override
@@ -154,12 +154,12 @@ struct PropertiesArray : public BasicMutableCommandBufferTest
    cl_mutable_command_khr command = nullptr;
 };

-struct Kernel : public BasicMutableCommandBufferTest
+struct Kernel : public InfoMutableCommandBufferTest
 {
-    using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest;
+    using InfoMutableCommandBufferTest::InfoMutableCommandBufferTest;

    Kernel(cl_device_id device, cl_context context, cl_command_queue queue)
-        : BasicMutableCommandBufferTest(device, context, queue)
+        : InfoMutableCommandBufferTest(device, context, queue)
    {}

    cl_int Run() override
@@ -195,12 +195,12 @@ struct Kernel : public BasicMutableCommandBufferTest
    cl_mutable_command_khr command = nullptr;
 };

-struct Dimensions : public BasicMutableCommandBufferTest
+struct Dimensions : public InfoMutableCommandBufferTest
 {
-    using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest;
+    using InfoMutableCommandBufferTest::InfoMutableCommandBufferTest;

    Dimensions(cl_device_id device, cl_context context, cl_command_queue queue)
-        : BasicMutableCommandBufferTest(device, context, queue)
+        : InfoMutableCommandBufferTest(device, context, queue)
    {}

    cl_int Run() override
@@ -234,12 +234,12 @@ struct Dimensions : public BasicMutableCommandBufferTest
    const size_t dimensions = 3;
 };

-struct InfoType : public BasicMutableCommandBufferTest
+struct InfoType : public InfoMutableCommandBufferTest
 {
-    using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest;
+    using InfoMutableCommandBufferTest::InfoMutableCommandBufferTest;

    InfoType(cl_device_id device, cl_context context, cl_command_queue queue)
-        : BasicMutableCommandBufferTest(device, context, queue)
+        : InfoMutableCommandBufferTest(device, context, queue)
    {}

    cl_int Run() override
@@ -271,12 +271,12 @@ struct InfoType : public BasicMutableCommandBufferTest
    cl_mutable_command_khr command = nullptr;
 };

-struct InfoQueue : public BasicMutableCommandBufferTest
+struct InfoQueue : public InfoMutableCommandBufferTest
 {
-    using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest;
+    using InfoMutableCommandBufferTest::InfoMutableCommandBufferTest;

    InfoQueue(cl_device_id device, cl_context context, cl_command_queue queue)
-        : BasicMutableCommandBufferTest(device, context, queue)
+        : InfoMutableCommandBufferTest(device, context, queue)
    {}

    cl_int Run() override
@@ -308,13 +308,13 @@ struct InfoQueue : public BasicMutableCommandBufferTest
    cl_mutable_command_khr command = nullptr;
 };

-struct InfoGlobalWorkOffset : public BasicMutableCommandBufferTest
+struct InfoGlobalWorkOffset : public InfoMutableCommandBufferTest
 {
-    using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest;
+    using InfoMutableCommandBufferTest::InfoMutableCommandBufferTest;

    InfoGlobalWorkOffset(cl_device_id device, cl_context context,
                         cl_command_queue queue)
-        : BasicMutableCommandBufferTest(device, context, queue)
+        : InfoMutableCommandBufferTest(device, context, queue)
    {}

    cl_int Run() override
@@ -346,13 +346,13 @@ struct InfoGlobalWorkOffset : public BasicMutableCommandBufferTest
    size_t test_global_work_offset = 0;
 };

-struct InfoGlobalWorkSize : public BasicMutableCommandBufferTest
+struct InfoGlobalWorkSize : public InfoMutableCommandBufferTest
 {
-    using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest;
+    using InfoMutableCommandBufferTest::InfoMutableCommandBufferTest;

    InfoGlobalWorkSize(cl_device_id device, cl_context context,
                       cl_command_queue queue)
-        : BasicMutableCommandBufferTest(device, context, queue)
+        : InfoMutableCommandBufferTest(device, context, queue)
    {}

    cl_int Run() override
@@ -383,13 +383,13 @@ struct InfoGlobalWorkSize : public BasicMutableCommandBufferTest
    size_t test_global_work_size = 0;
 };

-struct InfoLocalWorkSize : public BasicMutableCommandBufferTest
+struct InfoLocalWorkSize : public InfoMutableCommandBufferTest
 {
-    using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest;
+    using InfoMutableCommandBufferTest::InfoMutableCommandBufferTest;

    InfoLocalWorkSize(cl_device_id device, cl_context context,
                      cl_command_queue queue)
-        : BasicMutableCommandBufferTest(device, context, queue)
+        : InfoMutableCommandBufferTest(device, context, queue)
    {}

    cl_int Run() override
--- 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
@@ -59,4 +59,8 @@ extern int test_mutable_command_info_global_work_size(cl_device_id device,
                                                      cl_context context,
                                                      cl_command_queue queue,
                                                      int num_elements);
-#endif // CL_KHR_COMMAND_BUFFER_MUTABLE_DISPATCH_PROCS_H
+extern int test_mutable_dispatch_global_offset(cl_device_id device,
+                                               cl_context context,
+                                               cl_command_queue queue,
+                                               int num_elements);
+#endif /*_CL_KHR_COMMAND_BUFFER_MUTABLE_DISPATCH_PROCS_H*/