From 56974a58585b8c66d9beddccd984990e45ca0ad7 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Jastrz=C4=99bski?= Date: Tue, 27 Jun 2023 17:54:14 +0200 Subject: [PATCH] Add global offset tests for cl_khr_command_buffer_mutable_dispatch. (#1743) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * Add global offset tests for cl_khr_command_buffer_mutable_dispatch. Signed-off-by: Paweł Jastrzębski * 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 * 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 * Fix review comments. Changes made: - Fix skip conditions - Remove obsolete variable - Replace a variable with a constant Signed-off-by: Paweł Jastrzębski * Fix review comments. Changes made: - Remove explicit base class call - Fix constant magic number Signed-off-by: Paweł Jastrzębski * Fix constant magic number. Signed-off-by: Paweł Jastrzębski * Fix clang-format. Signed-off-by: Paweł Jastrzębski * Fix condition for result check. Signed-off-by: Paweł Jastrzębski --------- Signed-off-by: Paweł Jastrzębski --- .../CMakeLists.txt | 1 + .../main.cpp | 1 + .../mutable_command_basic.h | 63 ++++-- .../mutable_command_global_offset.cpp | 179 ++++++++++++++++++ .../mutable_command_info.cpp | 60 +++--- .../procs.h | 6 +- 6 files changed, 267 insertions(+), 43 deletions(-) create mode 100644 test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_global_offset.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 e0625833..80214609 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 @@ -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 ) 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 97075792..b53914dc 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 @@ -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[]) 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 96669583..c88c14d1 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 @@ -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( \ + 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( \ - 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 diff --git 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 new file mode 100644 index 00000000..70e1d9b1 --- /dev/null +++ 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 +#include "typeWrappers.h" +#include "procs.h" +#include "testHarness.h" +#include "imageHelpers.h" +#include +#include +#include +#include +#include +#include +#include "mutable_command_basic.h" + +#include +#include + +//////////////////////////////////////////////////////////////////////////////// +// 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 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(device, context, queue, + num_elements); +} diff --git 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 index cc425a4d..a8ed325a 100644 --- 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 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 4b6dacb6..588bdc81 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 @@ -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*/