From 37b8ecfb0ec58bb01493bb5a5887fd380c3d8a77 Mon Sep 17 00:00:00 2001 From: Michal Babej Date: Thu, 17 Oct 2024 18:16:29 +0300 Subject: [PATCH 1/3] cl_khr_command_buffer_mutable_dispatch: fix incorrect logic in Skip check the test checks 'cl_khr_extended_versioning' then immediately calls: get_extension_version(device, "cl_khr_command_buffer_mutable_dispatch"); .. however, get_extension_version throws an exception when called with unsupported extension. --- .../mutable_command_info.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) 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 2b5b3b42e8..d9edccd492 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 @@ -120,6 +120,7 @@ struct PropertiesArray : public InfoMutableCommandBufferTest virtual bool Skip() override { Version device_version = get_device_cl_version(device); + if (InfoMutableCommandBufferTest::Skip()) return true; if ((device_version >= Version(3, 0)) || is_extension_available(device, "cl_khr_extended_versioning")) { @@ -132,7 +133,7 @@ struct PropertiesArray : public InfoMutableCommandBufferTest return true; } } - return InfoMutableCommandBufferTest::Skip(); + return false; } cl_int Run() override From 7484f2d31a90b92539dc097135dc085889676c04 Mon Sep 17 00:00:00 2001 From: Michal Babej <90404+franz@users.noreply.github.com> Date: Wed, 6 Nov 2024 20:17:15 +0200 Subject: [PATCH 2/3] mutable_command_info: fix Skip logic in PropertiesArray test the test uses CL_MUTABLE_DISPATCH_ARGUMENTS_KHR, but does not check for support. --- .../mutable_command_info.cpp | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) 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 d9edccd492..4dd7d66efb 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 @@ -124,7 +124,6 @@ struct PropertiesArray : public InfoMutableCommandBufferTest if ((device_version >= Version(3, 0)) || is_extension_available(device, "cl_khr_extended_versioning")) { - cl_version extension_version = get_extension_version( device, "cl_khr_command_buffer_mutable_dispatch"); @@ -133,6 +132,16 @@ struct PropertiesArray : public InfoMutableCommandBufferTest return true; } } + + 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 & CL_MUTABLE_DISPATCH_ARGUMENTS_KHR) == 0) + return true; + return false; } From a2fc5d8d1c5d6191a4fdbf5e5a0df8064255b648 Mon Sep 17 00:00:00 2001 From: Michal Babej Date: Mon, 28 Oct 2024 17:48:13 +0200 Subject: [PATCH 3/3] fix workgroup sizes in mutable_command_global_size test test MutableDispatchGlobalSize was failing (with an implementation that doesn't support non-uniform WGs), because the update_global_size = 3 was not a multiple of the local work-size. Fixed by increasing the global work-size to 256K and update size to 16K. This should work with all devices that have max_work_group_size <= 16K. Updates also MutableDispatchWorkGroups which had an out-of-bounds access, because it hardcoded the value of global work-size. --- .../mutable_command_basic.h | 2 +- .../mutable_command_global_size.cpp | 6 +++--- .../mutable_command_work_groups.cpp | 5 ++--- 3 files changed, 6 insertions(+), 7 deletions(-) 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 51938dce02..c1697accbd 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 @@ -126,7 +126,7 @@ struct BasicMutableCommandBufferTest : BasicCommandBufferTest clUpdateMutableCommandsKHR_fn clUpdateMutableCommandsKHR = nullptr; const char* kernelString = "__kernel void empty() {}"; - const size_t global_work_size = 4 * 16; + const size_t global_work_size = 256 * 1024; }; struct InfoMutableCommandBufferTest : BasicMutableCommandBufferTest diff --git a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_global_size.cpp b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_global_size.cpp index 946fa995b3..22a58ff593 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_global_size.cpp +++ b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_global_size.cpp @@ -153,9 +153,9 @@ struct MutableDispatchGlobalSize : public InfoMutableCommandBufferTest } size_t info_global_size = 0; - const size_t update_global_size = 3; - const size_t sizeToAllocate = global_work_size; - const size_t num_elements = sizeToAllocate / sizeof(cl_int); + const size_t update_global_size = 16 * 1024; + const size_t sizeToAllocate = global_work_size * sizeof(cl_int); + const size_t num_elements = global_work_size; cl_mutable_command_khr command = nullptr; }; diff --git a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_work_groups.cpp b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_work_groups.cpp index ad20fbe3b2..03b42e0df2 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_work_groups.cpp +++ b/test_conformance/extensions/cl_khr_command_buffer/cl_khr_command_buffer_mutable_dispatch/mutable_command_work_groups.cpp @@ -278,10 +278,9 @@ struct MutableDispatchWorkGroups : public BasicMutableCommandBufferTest Configuration config; size_t info_global_size = 0; - static constexpr size_t test_global_work_size = 64; - static constexpr size_t update_global_size = 16; + static constexpr size_t update_global_size = 16 * 1024; const size_t local_work_size = 8; - const size_t sizeToAllocate = 64 * sizeof(cl_int); + const size_t sizeToAllocate = global_work_size * sizeof(cl_int); }; int test_command_buffer_with_no_additional_work_groups(cl_device_id device,