From 3478f3d30466780282054a6bb140dc459e079256 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pawe=C5=82=20Jastrz=C4=99bski?= Date: Tue, 21 Mar 2023 17:12:34 +0100 Subject: [PATCH] Add mutable dispatch tests. (#1651) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * Add mutable dispatch tests. Signed-off-by: Paweł Jastrzębski * Mutable dispatch buffer tests redesigned to inherit from command buffer tests. Signed-off-by: Paweł Jastrzębski * Add remaining tests for clGetMutableCommandInfoKHR. Added tests for missing queries: * CL_MUTABLE_DISPATCH_PROPERTIES_ARRAY_KHR * CL_MUTABLE_DISPATCH_KERNEL_KHR * CL_MUTABLE_DISPATCH_DIMENSIONS_KHR Signed-off-by: Paweł Jastrzębski * Minor code cleanup. Signed-off-by: Paweł Jastrzębski * Introduce review changes. Introduce changes requested in review: - Remove bitwise chceck for CL_DEVICE_MUTABLE_DISPATCH_CAPABILITIES_KHR - Add error check for clFinalizeCommandBufferKHR - Add global_work_size to clCommandNDRangeKernelKHR for CL_MUTABLE_COMMAND_COMMAND_BUFFER_KHR - Move BasicMutableCommandBufferTest to a separate header file - Change name of file command_buffer_test_mutable_dispatch.cpp to mutable_command_info.cpp Signed-off-by: Paweł Jastrzębski * Set global_work_size on every use of clCommandNDRangeKernelKHR. Signed-off-by: Paweł Jastrzębski * Apply changes for review. - Add error check for init_extension_functions() - Check mutable_capabilities for non-zero - Replace clKernelWrapper with cl_kernel for CL_MUTABLE_DISPATCH_KERNEL_KHR test - Replace clCommandBufferWrapper with cl_command_buffer_khr for CL_MUTABLE_COMMAND_COMMAND_BUFFER_KHR test - Remove unneded test_command_buffer declarations - Check type for CL_MUTABLE_COMMAND_COMMAND_TYPE_KHR - Remove retain() from operator= in clCommandBufferWrapper Signed-off-by: Paweł Jastrzębski * Apply changes for review. Apply changes for review: - Fix header guards - Add copyright header - Change checks for test_kernel, test_queue Signed-off-by: Paweł Jastrzębski --------- Signed-off-by: Paweł Jastrzębski --- .../cl_khr_command_buffer/CMakeLists.txt | 2 + .../CMakeLists.txt | 9 + .../main.cpp | 42 ++ .../mutable_command_basic.h | 107 ++++ .../mutable_command_info.cpp | 497 ++++++++++++++++++ .../procs.h | 62 +++ 6 files changed, 719 insertions(+) create mode 100644 test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/CMakeLists.txt create mode 100644 test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/main.cpp create mode 100644 test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_basic.h create mode 100644 test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_info.cpp create mode 100644 test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/procs.h diff --git a/test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt b/test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt index 98b9eb7f..a09d51c6 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt +++ b/test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt @@ -15,3 +15,5 @@ set(${MODULE_NAME}_SOURCES ) include(../../CMakeCommon.txt) + +add_subdirectory( cl_khr_command_buffer_mutable_dispatch ) 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 new file mode 100644 index 00000000..e0625833 --- /dev/null +++ b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/CMakeLists.txt @@ -0,0 +1,9 @@ +set(MODULE_NAME CL_KHR_MUTABLE_DISPATCH) + +set(${MODULE_NAME}_SOURCES + main.cpp + mutable_command_info.cpp + ../basic_command_buffer.cpp +) + +include(../../../CMakeCommon.txt) 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 new file mode 100644 index 00000000..97075792 --- /dev/null +++ b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/main.cpp @@ -0,0 +1,42 @@ +// 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 "procs.h" +#include "harness/testHarness.h" + +test_definition test_list[] = { + ADD_TEST(mutable_command_info_device_query), + ADD_TEST(mutable_command_info_buffer), + ADD_TEST(mutable_command_properties_array), + ADD_TEST(mutable_command_kernel), + ADD_TEST(mutable_command_dimensions), + ADD_TEST(mutable_command_info_type), + ADD_TEST(mutable_command_info_queue), + ADD_TEST(mutable_command_info_global_work_offset), + ADD_TEST(mutable_command_info_local_work_size), + ADD_TEST(mutable_command_info_global_work_size), +}; + +int main(int argc, const char *argv[]) +{ + // A device may report the required properties of a queue that + // is compatible with command-buffers via the query + // CL_DEVICE_COMMAND_BUFFER_REQUIRED_QUEUE_PROPERTIES_KHR. We account + // for this in the tests themselves, rather than here, where we have a + // device to query. + const cl_command_queue_properties queue_properties = 0; + return runTestHarnessWithCheck(argc, argv, ARRAY_SIZE(test_list), test_list, + false, queue_properties, nullptr); + return 0; +} 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 new file mode 100644 index 00000000..9056a00d --- /dev/null +++ b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_basic.h @@ -0,0 +1,107 @@ +// +// Copyright (c) 2023 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. + +#ifndef _CL_KHR_MUTABLE_COMMAND_BASIC_H +#define _CL_KHR_MUTABLE_COMMAND_BASIC_H + +#include "../basic_command_buffer.h" +#include "../command_buffer_test_base.h" + +struct BasicMutableCommandBufferTest : BasicCommandBufferTest +{ + BasicMutableCommandBufferTest(cl_device_id device, cl_context context, + cl_command_queue queue) + : BasicCommandBufferTest(device, context, queue) + {} + + virtual cl_int SetUp(int elements) override + { + BasicCommandBufferTest::SetUp(elements); + + cl_int error = init_extension_functions(); + test_error(error, "Unable to initialise extension functions"); + + const cl_command_buffer_properties_khr props[] = { + CL_COMMAND_BUFFER_FLAGS_KHR, + CL_COMMAND_BUFFER_MUTABLE_KHR, + 0, + }; + + command_buffer = clCreateCommandBufferKHR(1, &queue, props, &error); + test_error(error, "Unable to create command buffer"); + + clProgramWrapper program = clCreateProgramWithSource( + context, 1, &kernelString, nullptr, &error); + test_error(error, "Unable to create program"); + + error = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr); + test_error(error, "Unable to build program"); + + kernel = clCreateKernel(program, "empty", &error); + test_error(error, "Unable to create kernel"); + + return error; + } + + bool Skip() override + { + bool extension_avaliable = + is_extension_available(device, + "cl_khr_command_buffer_mutable_dispatch") + == 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 != 0; + + return !mutable_support || !extension_avaliable + || BasicCommandBufferTest::Skip(); + } + + 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"); + + // 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(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 \ No newline at end of file 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 new file mode 100644 index 00000000..cc425a4d --- /dev/null +++ b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_info.cpp @@ -0,0 +1,497 @@ +// +// 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 +#include +#include +#include +#include +#include +#include "mutable_command_basic.h" + +#include +#include +//////////////////////////////////////////////////////////////////////////////// +// mutable dispatch tests which handle following cases: +// +// CL_DEVICE_MUTABLE_DISPATCH_CAPABILITIES_KHR +// CL_MUTABLE_COMMAND_COMMAND_QUEUE_KHR +// CL_MUTABLE_COMMAND_COMMAND_BUFFER_KHR +// CL_MUTABLE_DISPATCH_PROPERTIES_ARRAY_KHR +// CL_MUTABLE_DISPATCH_KERNEL_KHR +// CL_MUTABLE_DISPATCH_DIMENSIONS_KHR +// CL_MUTABLE_DISPATCH_GLOBAL_WORK_OFFSET_KHR +// CL_MUTABLE_DISPATCH_GLOBAL_WORK_SIZE_KHR +// CL_MUTABLE_DISPATCH_LOCAL_WORK_SIZE_KHR +// CL_MUTABLE_COMMAND_COMMAND_TYPE_KHR + +struct InfoDeviceQuery : public BasicMutableCommandBufferTest +{ + using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest; + + InfoDeviceQuery(cl_device_id device, cl_context context, + cl_command_queue queue) + : BasicMutableCommandBufferTest(device, context, queue) + {} + + cl_int Run() override + { + cl_mutable_dispatch_fields_khr mutable_capabilities; + + cl_int error = clGetDeviceInfo( + device, CL_DEVICE_MUTABLE_DISPATCH_CAPABILITIES_KHR, + sizeof(mutable_capabilities), &mutable_capabilities, nullptr); + test_error(error, "clGetDeviceInfo failed"); + + if (!mutable_capabilities) + { + log_error("Device does not support update arguments to a " + "mutable-dispatch."); + return TEST_FAIL; + } + + return CL_SUCCESS; + } +}; + +struct InfoBuffer : public BasicMutableCommandBufferTest +{ + using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest; + + InfoBuffer(cl_device_id device, cl_context context, cl_command_queue queue) + : BasicMutableCommandBufferTest(device, context, queue) + {} + + cl_int Run() override + { + cl_int error = clCommandNDRangeKernelKHR( + command_buffer, nullptr, nullptr, kernel, 1, nullptr, + &global_work_size, nullptr, 0, nullptr, nullptr, &command); + test_error(error, "clCommandNDRangeKernelKHR failed"); + + error = clGetMutableCommandInfoKHR( + command, CL_MUTABLE_COMMAND_COMMAND_BUFFER_KHR, + sizeof(test_command_buffer), &test_command_buffer, nullptr); + test_error(error, "clGetMutableCommandInfoKHR failed"); + + if (test_command_buffer != command_buffer) + { + log_error("ERROR: Incorrect command buffer returned from " + "clGetMutableCommandInfoKHR."); + return TEST_FAIL; + } + + error = clFinalizeCommandBufferKHR(command_buffer); + test_error(error, "clFinalizeCommandBufferKHR failed"); + + return CL_SUCCESS; + } + + cl_command_buffer_khr test_command_buffer = nullptr; + cl_mutable_command_khr command = nullptr; +}; + +struct PropertiesArray : public BasicMutableCommandBufferTest +{ + using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest; + + PropertiesArray(cl_device_id device, cl_context context, + cl_command_queue queue) + : BasicMutableCommandBufferTest(device, context, queue) + {} + + cl_int Run() override + { + cl_ndrange_kernel_command_properties_khr props[] = { + CL_MUTABLE_DISPATCH_UPDATABLE_FIELDS_KHR, + CL_MUTABLE_DISPATCH_ARGUMENTS_KHR, 0 + }; + + cl_int error = clCommandNDRangeKernelKHR( + command_buffer, nullptr, props, kernel, 1, nullptr, + &global_work_size, nullptr, 0, nullptr, nullptr, &command); + test_error(error, "clCommandNDRangeKernelKHR failed"); + + cl_ndrange_kernel_command_properties_khr test_props[] = { 0, 0, 0 }; + size_t size; + + error = clGetMutableCommandInfoKHR( + command, CL_MUTABLE_DISPATCH_PROPERTIES_ARRAY_KHR, + sizeof(test_props), test_props, &size); + test_error(error, "clGetMutableCommandInfoKHR failed"); + + if (size != sizeof(props) || test_props[0] != props[0] + || test_props[1] != props[1]) + { + log_error("ERROR: Incorrect command buffer returned from " + "clGetMutableCommandInfoKHR."); + return TEST_FAIL; + } + + error = clFinalizeCommandBufferKHR(command_buffer); + test_error(error, "clFinalizeCommandBufferKHR failed"); + + return CL_SUCCESS; + } + + cl_mutable_command_khr command = nullptr; +}; + +struct Kernel : public BasicMutableCommandBufferTest +{ + using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest; + + Kernel(cl_device_id device, cl_context context, cl_command_queue queue) + : BasicMutableCommandBufferTest(device, context, queue) + {} + + cl_int Run() override + { + cl_int error = clCommandNDRangeKernelKHR( + command_buffer, nullptr, nullptr, kernel, 1, nullptr, + &global_work_size, nullptr, 0, nullptr, nullptr, &command); + test_error(error, "clCommandNDRangeKernelKHR failed"); + + cl_kernel test_kernel; + size_t size; + + error = clGetMutableCommandInfoKHR( + command, CL_MUTABLE_DISPATCH_KERNEL_KHR, sizeof(test_kernel), + &test_kernel, &size); + test_error(error, "clGetMutableCommandInfoKHR failed"); + + // We can not check if this is the right kernel because this is an + // opaque object. + if (test_kernel != kernel) + { + log_error("ERROR: Incorrect command buffer returned from " + "clGetMutableCommandInfoKHR."); + return TEST_FAIL; + } + + error = clFinalizeCommandBufferKHR(command_buffer); + test_error(error, "clFinalizeCommandBufferKHR failed"); + + return CL_SUCCESS; + } + + cl_mutable_command_khr command = nullptr; +}; + +struct Dimensions : public BasicMutableCommandBufferTest +{ + using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest; + + Dimensions(cl_device_id device, cl_context context, cl_command_queue queue) + : BasicMutableCommandBufferTest(device, context, queue) + {} + + cl_int Run() override + { + cl_int error = clCommandNDRangeKernelKHR( + command_buffer, nullptr, nullptr, kernel, dimensions, nullptr, + &global_work_size, nullptr, 0, nullptr, nullptr, &command); + test_error(error, "clCommandNDRangeKernelKHR failed"); + + size_t test_dimensions; + + error = clGetMutableCommandInfoKHR( + command, CL_MUTABLE_DISPATCH_DIMENSIONS_KHR, + sizeof(test_dimensions), &test_dimensions, nullptr); + test_error(error, "clGetMutableCommandInfoKHR failed"); + + if (test_dimensions != dimensions) + { + log_error("ERROR: Incorrect command buffer returned from " + "clGetMutableCommandInfoKHR."); + return TEST_FAIL; + } + + error = clFinalizeCommandBufferKHR(command_buffer); + test_error(error, "clFinalizeCommandBufferKHR failed"); + + return CL_SUCCESS; + } + + cl_mutable_command_khr command = nullptr; + const size_t dimensions = 3; +}; + +struct InfoType : public BasicMutableCommandBufferTest +{ + using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest; + + InfoType(cl_device_id device, cl_context context, cl_command_queue queue) + : BasicMutableCommandBufferTest(device, context, queue) + {} + + cl_int Run() override + { + cl_int error = clCommandNDRangeKernelKHR( + command_buffer, nullptr, nullptr, kernel, 1, nullptr, + &global_work_size, nullptr, 0, nullptr, nullptr, &command); + test_error(error, "clCommandNDRangeKernelKHR failed"); + + cl_command_type type = 0; + error = clGetMutableCommandInfoKHR(command, + CL_MUTABLE_COMMAND_COMMAND_TYPE_KHR, + sizeof(type), &type, NULL); + test_error(error, "clGetMutableCommandInfoKHR failed"); + + if (type != CL_COMMAND_NDRANGE_KERNEL) + { + log_error("ERROR: Wrong type returned from " + "clGetMutableCommandInfoKHR."); + return TEST_FAIL; + } + + error = clFinalizeCommandBufferKHR(command_buffer); + test_error(error, "clFinalizeCommandBufferKHR failed"); + + return CL_SUCCESS; + } + + cl_mutable_command_khr command = nullptr; +}; + +struct InfoQueue : public BasicMutableCommandBufferTest +{ + using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest; + + InfoQueue(cl_device_id device, cl_context context, cl_command_queue queue) + : BasicMutableCommandBufferTest(device, context, queue) + {} + + cl_int Run() override + { + cl_int error = clCommandNDRangeKernelKHR( + command_buffer, nullptr, nullptr, kernel, 1, nullptr, + &global_work_size, nullptr, 0, nullptr, nullptr, &command); + test_error(error, "clCommandNDRangeKernelKHR failed"); + + cl_command_queue test_queue = nullptr; + error = clGetMutableCommandInfoKHR( + command, CL_MUTABLE_COMMAND_COMMAND_QUEUE_KHR, sizeof(test_queue), + &test_queue, nullptr); + test_error(error, "clGetMutableCommandInfoKHR failed"); + + if (test_queue != queue) + { + log_error("ERROR: Incorrect queue returned from " + "clGetMutableCommandInfoKHR."); + return TEST_FAIL; + } + + error = clFinalizeCommandBufferKHR(command_buffer); + test_error(error, "clFinalizeCommandBufferKHR failed"); + + return CL_SUCCESS; + } + + cl_mutable_command_khr command = nullptr; +}; + +struct InfoGlobalWorkOffset : public BasicMutableCommandBufferTest +{ + using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest; + + InfoGlobalWorkOffset(cl_device_id device, cl_context context, + cl_command_queue queue) + : BasicMutableCommandBufferTest(device, context, queue) + {} + + cl_int Run() override + { + cl_int error = clCommandNDRangeKernelKHR( + command_buffer, nullptr, nullptr, kernel, 1, &global_work_offset, + &global_work_size, nullptr, 0, nullptr, nullptr, &command); + test_error(error, "clCommandNDRangeKernelKHR failed"); + + error = clGetMutableCommandInfoKHR( + command, CL_MUTABLE_DISPATCH_GLOBAL_WORK_OFFSET_KHR, + sizeof(test_global_work_offset), &test_global_work_offset, nullptr); + + if (test_global_work_offset != global_work_offset) + { + log_error("ERROR: Wrong size returned from " + "clGetMutableCommandInfoKHR."); + return TEST_FAIL; + } + + error = clFinalizeCommandBufferKHR(command_buffer); + test_error(error, "clFinalizeCommandBufferKHR failed"); + + return CL_SUCCESS; + } + + cl_mutable_command_khr command = nullptr; + const size_t global_work_offset = 4 * sizeof(cl_int); + size_t test_global_work_offset = 0; +}; + +struct InfoGlobalWorkSize : public BasicMutableCommandBufferTest +{ + using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest; + + InfoGlobalWorkSize(cl_device_id device, cl_context context, + cl_command_queue queue) + : BasicMutableCommandBufferTest(device, context, queue) + {} + + cl_int Run() override + { + cl_int error = clCommandNDRangeKernelKHR( + command_buffer, nullptr, nullptr, kernel, 1, nullptr, + &global_work_size, nullptr, 0, nullptr, nullptr, &command); + test_error(error, "clCommandNDRangeKernelKHR failed"); + + error = clGetMutableCommandInfoKHR( + command, CL_MUTABLE_DISPATCH_GLOBAL_WORK_SIZE_KHR, + sizeof(test_global_work_size), &test_global_work_size, nullptr); + + if (test_global_work_size != global_work_size) + { + log_error("ERROR: Wrong size returned from " + "clGetMutableCommandInfoKHR."); + return TEST_FAIL; + } + + error = clFinalizeCommandBufferKHR(command_buffer); + test_error(error, "clFinalizeCommandBufferKHR failed"); + + return TEST_PASS; + } + + cl_mutable_command_khr command = nullptr; + size_t test_global_work_size = 0; +}; + +struct InfoLocalWorkSize : public BasicMutableCommandBufferTest +{ + using BasicMutableCommandBufferTest::BasicMutableCommandBufferTest; + + InfoLocalWorkSize(cl_device_id device, cl_context context, + cl_command_queue queue) + : BasicMutableCommandBufferTest(device, context, queue) + {} + + cl_int Run() override + { + cl_int 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 = clGetMutableCommandInfoKHR( + command, CL_MUTABLE_DISPATCH_LOCAL_WORK_SIZE_KHR, + sizeof(test_local_work_size), &test_local_work_size, nullptr); + + if (test_local_work_size != local_work_size) + { + log_error("ERROR: Wrong size returned from " + "clGetMutableCommandInfoKHR."); + return TEST_FAIL; + } + + error = clFinalizeCommandBufferKHR(command_buffer); + test_error(error, "clFinalizeCommandBufferKHR failed"); + + return CL_SUCCESS; + } + + cl_mutable_command_khr command = nullptr; + const size_t local_work_size = 4 * sizeof(cl_int); + size_t test_local_work_size = 0; +}; + +int test_mutable_command_info_device_query(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements) +{ + return MakeAndRunTest(device, context, queue, + num_elements); +} + +int test_mutable_command_info_buffer(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return MakeAndRunTest(device, context, queue, num_elements); +} + +int test_mutable_command_properties_array(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements) +{ + return MakeAndRunTest(device, context, queue, + num_elements); +} + +int test_mutable_command_kernel(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return MakeAndRunTest(device, context, queue, num_elements); +} + +int test_mutable_command_dimensions(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return MakeAndRunTest(device, context, queue, num_elements); +} + +int test_mutable_command_info_type(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return MakeAndRunTest(device, context, queue, num_elements); +} + +int test_mutable_command_info_queue(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return MakeAndRunTest(device, context, queue, num_elements); +} + +int test_mutable_command_info_global_work_offset(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements) +{ + return MakeAndRunTest(device, context, queue, + num_elements); +} + +int test_mutable_command_info_global_work_size(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements) +{ + return MakeAndRunTest(device, context, queue, + num_elements); +} + +int test_mutable_command_info_local_work_size(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 new file mode 100644 index 00000000..08512cae --- /dev/null +++ b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/procs.h @@ -0,0 +1,62 @@ +// +// Copyright (c) 2023 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. +// +#ifndef _CL_KHR_COMMAND_BUFFER_MUTABLE_DISPATCH_PROCS_H +#define _CL_KHR_COMMAND_BUFFER_MUTABLE_DISPATCH_PROCS_H + +#include + + +// Basic mutable dispatch tests +extern int test_mutable_command_info_device_query(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_mutable_command_info_buffer(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_mutable_command_info_type(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_mutable_command_info_queue(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_mutable_command_properties_array(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_mutable_command_kernel(cl_device_id device, cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_mutable_command_dimensions(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_mutable_command_info_global_work_offset(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_mutable_command_info_local_work_size(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements); +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*/