From 1139f54bde0e8c9dc7da1d03769ec8d7dfe42310 Mon Sep 17 00:00:00 2001 From: Sven van Haastregt Date: Mon, 15 Apr 2024 13:37:07 +0200 Subject: [PATCH 1/7] device_execution: fix -Wformat warnings (#1938) Printing of a `size_t` requires the `%zu` specifier. Signed-off-by: Sven van Haastregt --- .../device_execution/enqueue_block.cpp | 7 ++- .../device_execution/enqueue_flags.cpp | 7 ++- .../device_execution/enqueue_ndrange.cpp | 7 ++- .../device_execution/enqueue_wg_size.cpp | 61 +++++++++++++------ .../device_execution/execute_block.cpp | 7 ++- .../device_execution/nested_blocks.cpp | 7 ++- 6 files changed, 66 insertions(+), 30 deletions(-) diff --git a/test_conformance/device_execution/enqueue_block.cpp b/test_conformance/device_execution/enqueue_block.cpp index b3cd154f9..46407b191 100644 --- a/test_conformance/device_execution/enqueue_block.cpp +++ b/test_conformance/device_execution/enqueue_block.cpp @@ -641,7 +641,9 @@ int test_enqueue_block(cl_device_id device, cl_context context, cl_command_queue if (!gKernelName.empty() && gKernelName != sources_enqueue_block[i].kernel_name) continue; - log_info("Running '%s' kernel (%d of %d) ...\n", sources_enqueue_block[i].kernel_name, i + 1, num_kernels_enqueue_block); + log_info("Running '%s' kernel (%d of %zu) ...\n", + sources_enqueue_block[i].kernel_name, i + 1, + num_kernels_enqueue_block); err_ret = run_n_kernel_args(context, queue, sources_enqueue_block[i].lines, sources_enqueue_block[i].num_lines, sources_enqueue_block[i].kernel_name, local_size, global_size, kernel_results, sizeof(kernel_results), 0, NULL); if(check_error(err_ret, "'%s' kernel execution failed", sources_enqueue_block[i].kernel_name)) { ++failCnt; res = -1; } else if((n = check_kernel_results(kernel_results, arr_size(kernel_results))) >= 0 && check_error(-1, "'%s' kernel results validation failed: [%d] returned %d expected 0", sources_enqueue_block[i].kernel_name, n, kernel_results[n])) res = -1; @@ -650,7 +652,8 @@ int test_enqueue_block(cl_device_id device, cl_context context, cl_command_queue if (failCnt > 0) { - log_error("ERROR: %d of %d kernels failed.\n", failCnt, num_kernels_enqueue_block); + log_error("ERROR: %zu of %zu kernels failed.\n", failCnt, + num_kernels_enqueue_block); } return res; diff --git a/test_conformance/device_execution/enqueue_flags.cpp b/test_conformance/device_execution/enqueue_flags.cpp index 6ae1f281d..44902d7f4 100644 --- a/test_conformance/device_execution/enqueue_flags.cpp +++ b/test_conformance/device_execution/enqueue_flags.cpp @@ -714,7 +714,9 @@ int test_enqueue_flags(cl_device_id device, cl_context context, cl_command_queue if (!gKernelName.empty() && gKernelName != sources_enqueue_block_flags[i].kernel_name) continue; - log_info("Running '%s' kernel (%d of %d) ...\n", sources_enqueue_block_flags[i].kernel_name, i + 1, num_enqueue_block_flags); + log_info("Running '%s' kernel (%d of %zu) ...\n", + sources_enqueue_block_flags[i].kernel_name, i + 1, + num_enqueue_block_flags); clMemWrapper mem = clCreateBuffer(context, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, global_size * BITS_DEPTH * sizeof(cl_int), buff, &err_ret); test_error(err_ret, "clCreateBuffer() failed"); @@ -749,7 +751,8 @@ int test_enqueue_flags(cl_device_id device, cl_context context, cl_command_queue if (failCnt > 0) { - log_error("ERROR: %d of %d kernels failed.\n", failCnt, num_enqueue_block_flags); + log_error("ERROR: %zu of %zu kernels failed.\n", failCnt, + num_enqueue_block_flags); } return res; diff --git a/test_conformance/device_execution/enqueue_ndrange.cpp b/test_conformance/device_execution/enqueue_ndrange.cpp index 42abe78b1..bc28b5efe 100644 --- a/test_conformance/device_execution/enqueue_ndrange.cpp +++ b/test_conformance/device_execution/enqueue_ndrange.cpp @@ -700,7 +700,9 @@ int test_enqueue_ndrange(cl_device_id device, cl_context context, cl_command_que { sizeof(cl_mem), &mem4 }, }; - log_info("Running '%s' kernel (%d of %d) ...\n", sources_ndrange_Xd[i].src.kernel_name, i + 1, num_kernels_ndrange_Xd); + log_info("Running '%s' kernel (%d of %zu) ...\n", + sources_ndrange_Xd[i].src.kernel_name, i + 1, + num_kernels_ndrange_Xd); err_ret = run_single_kernel_args(context, queue, sources_ndrange_Xd[i].src.lines, sources_ndrange_Xd[i].src.num_lines, sources_ndrange_Xd[i].src.kernel_name, kernel_results, sizeof(kernel_results), arr_size(args), args); cl_int *ptr = (cl_int *)clEnqueueMapBuffer(queue, mem3, CL_TRUE, CL_MAP_READ, 0, glob_results.size() * sizeof(cl_int), 0, 0, 0, &err_ret); @@ -718,7 +720,8 @@ int test_enqueue_ndrange(cl_device_id device, cl_context context, cl_command_que if (failCnt > 0) { - log_error("ERROR: %d of %d kernels failed.\n", failCnt, num_kernels_ndrange_Xd); + log_error("ERROR: %zu of %zu kernels failed.\n", failCnt, + num_kernels_ndrange_Xd); } return res; diff --git a/test_conformance/device_execution/enqueue_wg_size.cpp b/test_conformance/device_execution/enqueue_wg_size.cpp index a037945c0..3da2d0a82 100644 --- a/test_conformance/device_execution/enqueue_wg_size.cpp +++ b/test_conformance/device_execution/enqueue_wg_size.cpp @@ -68,13 +68,15 @@ static int check_single(cl_int* results, cl_int len, cl_int nesting_level) { if(i == 0 && results[i] != nestingLevel) { - log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], nestingLevel, i); + log_error("ERROR: Kernel returned %d vs. expected %d, index: %zu\n", + results[i], nestingLevel, i); return (int)i; } if(i > 0 && results[i] != 0) { - log_error("ERROR: Kernel returned %d vs. expected 0, index: %d\n", results[i], i); + log_error("ERROR: Kernel returned %d vs. expected 0, index: %zu\n", + results[i], i); return (int)i; } } @@ -142,7 +144,8 @@ static int check_some_eq_1D(cl_int* results, cl_int len, cl_int nesting_level) { if (results[i] != referenceResults[i]) { - log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); + log_error("ERROR: Kernel returned %d vs. expected %d, index: %zu\n", + results[i], referenceResults[i], i); return (int)i; } } @@ -216,7 +219,8 @@ static int check_some_diff_1D(cl_int* results, cl_int maxGlobalWorkSize, cl_int { if (results[i] != referenceResults[i]) { - log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); + log_error("ERROR: Kernel returned %d vs. expected %d, index: %zu\n", + results[i], referenceResults[i], i); return (int)i; } } @@ -278,7 +282,8 @@ static int check_all_eq_1D(cl_int* results, cl_int len, cl_int nesting_level) { if (results[i] != referenceResults[i]) { - log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); + log_error("ERROR: Kernel returned %d vs. expected %d, index: %zu\n", + results[i], referenceResults[i], i); return (int)i; } } @@ -346,7 +351,8 @@ static int check_all_diff_1D(cl_int* results, cl_int len, cl_int nesting_level) { if (results[i] != referenceResults[i]) { - log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); + log_error("ERROR: Kernel returned %d vs. expected %d, index: %zu\n", + results[i], referenceResults[i], i); return (int)i; } } @@ -455,7 +461,8 @@ static int check_some_eq_2D(cl_int* results, cl_int len, cl_int nesting_level) { if (results[i] != referenceResults[i]) { - log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); + log_error("ERROR: Kernel returned %d vs. expected %d, index: %zu\n", + results[i], referenceResults[i], i); return (int)i; } } @@ -535,7 +542,8 @@ static int check_some_diff_2D(cl_int* results, cl_int len, cl_int nesting_level) { if (results[i] != referenceResults[i]) { - log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); + log_error("ERROR: Kernel returned %d vs. expected %d, index: %zu\n", + results[i], referenceResults[i], i); return (int)i; } } @@ -603,7 +611,8 @@ static int check_all_eq_2D(cl_int* results, cl_int len, cl_int nesting_level) { if (results[i] != referenceResults[i]) { - log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); + log_error("ERROR: Kernel returned %d vs. expected %d, index: %zu\n", + results[i], referenceResults[i], i); return (int)i; } } @@ -677,7 +686,8 @@ static int check_all_diff_2D(cl_int* results, cl_int len, cl_int nesting_level) { if (results[i] != referenceResults[i]) { - log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); + log_error("ERROR: Kernel returned %d vs. expected %d, index: %zu\n", + results[i], referenceResults[i], i); return (int)i; } } @@ -794,7 +804,8 @@ static int check_some_eq_3D(cl_int* results, cl_int len, cl_int nesting_level) { if (results[i] != referenceResults[i]) { - log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); + log_error("ERROR: Kernel returned %d vs. expected %d, index: %zu\n", + results[i], referenceResults[i], i); return (int)i; } } @@ -879,7 +890,8 @@ static int check_some_diff_3D(cl_int* results, cl_int len, cl_int nesting_level) { if (results[i] != referenceResults[i]) { - log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); + log_error("ERROR: Kernel returned %d vs. expected %d, index: %zu\n", + results[i], referenceResults[i], i); return (int)i; } } @@ -952,7 +964,8 @@ static int check_all_eq_3D(cl_int* results, cl_int len, cl_int nesting_level) { if (results[i] != referenceResults[i]) { - log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); + log_error("ERROR: Kernel returned %d vs. expected %d, index: %zu\n", + results[i], referenceResults[i], i); return (int)i; } } @@ -1031,7 +1044,8 @@ static int check_all_diff_3D(cl_int* results, cl_int len, cl_int nesting_level) { if (results[i] != referenceResults[i]) { - log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); + log_error("ERROR: Kernel returned %d vs. expected %d, index: %zu\n", + results[i], referenceResults[i], i); return (int)i; } } @@ -1217,7 +1231,8 @@ static int check_some_eq_mix(cl_int* results, cl_int len, cl_int nesting_level) { if (results[i] != referenceResults[i]) { - log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); + log_error("ERROR: Kernel returned %d vs. expected %d, index: %zu\n", + results[i], referenceResults[i], i); return (int)i; } } @@ -1346,7 +1361,8 @@ static int check_some_diff_mix(cl_int* results, cl_int len, cl_int nesting_level { if (results[i] != referenceResults[i]) { - log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); + log_error("ERROR: Kernel returned %d vs. expected %d, index: %zu\n", + results[i], referenceResults[i], i); return (int)i; } } @@ -1462,7 +1478,8 @@ static int check_all_eq_mix(cl_int* results, cl_int len, cl_int nesting_level) { if (results[i] != referenceResults[i]) { - log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); + log_error("ERROR: Kernel returned %d vs. expected %d, index: %zu\n", + results[i], referenceResults[i], i); return (int)i; } } @@ -1584,7 +1601,8 @@ static int check_all_diff_mix(cl_int* results, cl_int len, cl_int nesting_level) { if (results[i] != referenceResults[i]) { - log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); + log_error("ERROR: Kernel returned %d vs. expected %d, index: %zu\n", + results[i], referenceResults[i], i); return (int)i; } } @@ -1670,7 +1688,9 @@ int test_enqueue_wg_size(cl_device_id device, cl_context context, cl_command_que if (!gKernelName.empty() && gKernelName != sources_enqueue_wg_size[k].src.kernel_name) continue; - log_info("Running '%s' kernel (%d of %d) ...\n", sources_enqueue_wg_size[k].src.kernel_name, k + 1, arr_size(sources_enqueue_wg_size)); + log_info("Running '%s' kernel (%d of %zu) ...\n", + sources_enqueue_wg_size[k].src.kernel_name, k + 1, + arr_size(sources_enqueue_wg_size)); for(i = 0; i < MAX_GLOBAL_WORK_SIZE; ++i) { kernel_results[i] = 0; @@ -1714,7 +1734,8 @@ int test_enqueue_wg_size(cl_device_id device, cl_context context, cl_command_que if (failCnt > 0) { - log_error("ERROR: %d of %d kernels failed.\n", failCnt, arr_size(sources_enqueue_wg_size)); + log_error("ERROR: %zu of %zu kernels failed.\n", failCnt, + arr_size(sources_enqueue_wg_size)); } free_mtdata(d); diff --git a/test_conformance/device_execution/execute_block.cpp b/test_conformance/device_execution/execute_block.cpp index e5b13eff7..5b7ee2f1c 100644 --- a/test_conformance/device_execution/execute_block.cpp +++ b/test_conformance/device_execution/execute_block.cpp @@ -1031,7 +1031,9 @@ int test_execute_block(cl_device_id device, cl_context context, cl_command_queue if (!gKernelName.empty() && gKernelName != sources_execute_block[i].kernel_name) continue; - log_info("Running '%s' kernel (%d of %d) ...\n", sources_execute_block[i].kernel_name, i + 1, num_kernels_execute_block); + log_info("Running '%s' kernel (%zu of %zu) ...\n", + sources_execute_block[i].kernel_name, i + 1, + num_kernels_execute_block); err_ret = run_n_kernel_args(context, queue, sources_execute_block[i].lines, sources_execute_block[i].num_lines, sources_execute_block[i].kernel_name, local_size, global_size, kernel_results, sizeof(kernel_results), 0, NULL); if(check_error(err_ret, "'%s' kernel execution failed", sources_execute_block[i].kernel_name)) { ++failCnt; res = -1; } else if((n = check_kernel_results(kernel_results, arr_size(kernel_results))) >= 0 && check_error(-1, "'%s' kernel results validation failed: [%d] returned %d expected 0", sources_execute_block[i].kernel_name, n, kernel_results[n])) { ++failCnt; res = -1; } @@ -1040,7 +1042,8 @@ int test_execute_block(cl_device_id device, cl_context context, cl_command_queue if (failCnt > 0) { - log_error("ERROR: %d of %d kernels failed.\n", failCnt, num_kernels_execute_block); + log_error("ERROR: %zu of %zu kernels failed.\n", failCnt, + num_kernels_execute_block); } return res; diff --git a/test_conformance/device_execution/nested_blocks.cpp b/test_conformance/device_execution/nested_blocks.cpp index 3898a716e..3f41bee7d 100644 --- a/test_conformance/device_execution/nested_blocks.cpp +++ b/test_conformance/device_execution/nested_blocks.cpp @@ -350,7 +350,9 @@ int test_enqueue_nested_blocks(cl_device_id device, cl_context context, cl_comma if (!gKernelName.empty() && gKernelName != sources_nested_blocks[k].src.kernel_name) continue; - log_info("Running '%s' kernel (%d of %d) ...\n", sources_nested_blocks[k].src.kernel_name, k + 1, arr_size(sources_nested_blocks)); + log_info("Running '%s' kernel (%d of %zu) ...\n", + sources_nested_blocks[k].src.kernel_name, k + 1, + arr_size(sources_nested_blocks)); for(i = 0; i < MAX_GLOBAL_WORK_SIZE; ++i) kernel_results[i] = 0; err_ret = run_n_kernel_args(context, queue, sources_nested_blocks[k].src.lines, sources_nested_blocks[k].src.num_lines, sources_nested_blocks[k].src.kernel_name, 0, MAX_GLOBAL_WORK_SIZE, kernel_results, sizeof(kernel_results), arr_size(args), args); @@ -366,7 +368,8 @@ int test_enqueue_nested_blocks(cl_device_id device, cl_context context, cl_comma if (failCnt > 0) { - log_error("ERROR: %d of %d kernels failed.\n", failCnt, arr_size(sources_nested_blocks)); + log_error("ERROR: %zu of %zu kernels failed.\n", failCnt, + arr_size(sources_nested_blocks)); } return res; From 7fa567c7a5a9049c0fa47fd0dd8b1e155b18a4f3 Mon Sep 17 00:00:00 2001 From: Kamil-Goras-Mobica <141216953+kamil-goras-mobica@users.noreply.github.com> Date: Tue, 16 Apr 2024 17:40:44 +0200 Subject: [PATCH 2/7] Corrections for negative tests for function CreateCommandBufferKHR (#1915) --- .../cl_khr_command_buffer/CMakeLists.txt | 1 + .../extensions/cl_khr_command_buffer/main.cpp | 7 + .../negative_command_buffer_create.cpp | 322 ++++++++++++++++++ .../extensions/cl_khr_command_buffer/procs.h | 20 ++ 4 files changed, 350 insertions(+) create mode 100644 test_conformance/extensions/cl_khr_command_buffer/negative_command_buffer_create.cpp diff --git a/test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt b/test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt index 0159e2c2d..11f88bb35 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt +++ b/test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt @@ -16,6 +16,7 @@ set(${MODULE_NAME}_SOURCES command_buffer_test_barrier.cpp command_buffer_test_event_info.cpp command_buffer_finalize.cpp + negative_command_buffer_create.cpp negative_command_buffer_get_info.cpp ) diff --git a/test_conformance/extensions/cl_khr_command_buffer/main.cpp b/test_conformance/extensions/cl_khr_command_buffer/main.cpp index 8f9cdace9..bd6120ccc 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/main.cpp +++ b/test_conformance/extensions/cl_khr_command_buffer/main.cpp @@ -64,6 +64,13 @@ test_definition test_list[] = { ADD_TEST(event_info_reference_count), ADD_TEST(finalize_invalid), ADD_TEST(finalize_empty), + ADD_TEST(negative_create_command_buffer_num_queues), + ADD_TEST(negative_create_command_buffer_null_queues), + ADD_TEST(negative_create_command_buffer_repeated_properties), + ADD_TEST(negative_create_command_buffer_not_supported_properties), + ADD_TEST(negative_create_command_buffer_queue_without_min_properties), + ADD_TEST( + negative_create_command_buffer_device_does_not_support_out_of_order_queue), ADD_TEST(negative_get_command_buffer_info_invalid_command_buffer), ADD_TEST(negative_get_command_buffer_info_not_supported_param_name), ADD_TEST(negative_get_command_buffer_info_queues), diff --git a/test_conformance/extensions/cl_khr_command_buffer/negative_command_buffer_create.cpp b/test_conformance/extensions/cl_khr_command_buffer/negative_command_buffer_create.cpp new file mode 100644 index 000000000..5cc287c0b --- /dev/null +++ b/test_conformance/extensions/cl_khr_command_buffer/negative_command_buffer_create.cpp @@ -0,0 +1,322 @@ +// +// Copyright (c) 2024 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 "basic_command_buffer.h" +#include "procs.h" + + +//-------------------------------------------------------------------------- +namespace { + +// CL_INVALID_VALUE if num_queues is not one. +struct CreateCommandBufferNumQueues : public BasicCommandBufferTest +{ + CreateCommandBufferNumQueues(cl_device_id device, cl_context context, + cl_command_queue queue) + : BasicCommandBufferTest(device, context, queue), queue1(nullptr), + queue2(nullptr) + {} + + cl_int Run() override + { + cl_int error = CL_SUCCESS; + + cl_command_queue queues[2] = { queue1, queue2 }; + + command_buffer = clCreateCommandBufferKHR(2, queues, nullptr, &error); + + test_failure_error_ret( + error, CL_INVALID_VALUE, + "clCreateCommandBufferKHR should return CL_INVALID_VALUE", + TEST_FAIL); + + return CL_SUCCESS; + } + + cl_int SetUp(int elements) override + { + cl_int error = CL_SUCCESS; + + error = BasicCommandBufferTest::SetUp(elements); + test_error(error, "BasicCommandBufferTest::SetUp failed"); + + queue1 = clCreateCommandQueue(context, device, 0, &error); + test_error(error, "clCreateCommandQueue failed"); + queue2 = clCreateCommandQueue(context, device, 0, &error); + test_error(error, "clCreateCommandQueue failed"); + + return CL_SUCCESS; + } + + bool Skip() override + { + return BasicCommandBufferTest::Skip() + || is_extension_available(device, + "cl_khr_command_buffer_multi_device"); + } + + clCommandQueueWrapper queue1; + clCommandQueueWrapper queue2; +}; + +// CL_INVALID_VALUE if queues is NULL. +struct CreateCommandBufferNullQueues : public BasicCommandBufferTest +{ + using BasicCommandBufferTest::BasicCommandBufferTest; + + cl_int Run() override + { + cl_int error = CL_SUCCESS; + + command_buffer = clCreateCommandBufferKHR(1, nullptr, nullptr, &error); + + test_failure_error_ret( + error, CL_INVALID_VALUE, + "clCreateCommandBufferKHR should return CL_INVALID_VALUE", + TEST_FAIL); + + return CL_SUCCESS; + } +}; + +// CL_INVALID_VALUE if values specified in properties are not valid, +// or if the same property name is specified more than once. +struct CreateCommandBufferRepeatedProperties : public BasicCommandBufferTest +{ + using BasicCommandBufferTest::BasicCommandBufferTest; + + cl_int Run() override + { + cl_int error = CL_SUCCESS; + + cl_command_buffer_properties_khr repeated_properties[5] = { + CL_COMMAND_BUFFER_FLAGS_KHR, CL_COMMAND_BUFFER_SIMULTANEOUS_USE_KHR, + CL_COMMAND_BUFFER_FLAGS_KHR, CL_COMMAND_BUFFER_SIMULTANEOUS_USE_KHR, + 0 + }; + + command_buffer = + clCreateCommandBufferKHR(1, &queue, repeated_properties, &error); + test_failure_error_ret( + error, CL_INVALID_VALUE, + "clCreateCommandBufferKHR should return CL_INVALID_VALUE", + TEST_FAIL); + + cl_command_buffer_properties_khr invalid_properties[3] = { + CL_COMMAND_BUFFER_FLAGS_KHR, CL_INVALID_PROPERTY, 0 + }; + + command_buffer = + clCreateCommandBufferKHR(1, &queue, invalid_properties, &error); + test_failure_error_ret( + error, CL_INVALID_VALUE, + "clCreateCommandBufferKHR should return CL_INVALID_VALUE", + TEST_FAIL); + + return CL_SUCCESS; + } +}; + +// CL_INVALID_PROPERTY if values specified in properties are valid but are not +// supported by all the devices associated with command-queues in queues. +struct CreateCommandBufferNotSupportedProperties : public BasicCommandBufferTest +{ + using BasicCommandBufferTest::BasicCommandBufferTest; + + cl_int Run() override + { + cl_int error = CL_SUCCESS; + + cl_command_buffer_properties_khr properties[3] = { + CL_COMMAND_BUFFER_FLAGS_KHR, CL_COMMAND_BUFFER_SIMULTANEOUS_USE_KHR, + 0 + }; + + command_buffer = + clCreateCommandBufferKHR(1, &queue, properties, &error); + test_failure_error_ret( + error, CL_INVALID_PROPERTY, + "clCreateCommandBufferKHR should return CL_INVALID_PROPERTY", + TEST_FAIL); + + return CL_SUCCESS; + } + + bool Skip() override + { + cl_device_command_buffer_capabilities_khr capabilities; + cl_int error = + clGetDeviceInfo(device, CL_DEVICE_COMMAND_BUFFER_CAPABILITIES_KHR, + sizeof(capabilities), &capabilities, NULL); + test_error(error, + "Unable to query CL_DEVICE_COMMAND_BUFFER_CAPABILITIES_KHR"); + + bool device_supports_simultaneous_use = + (capabilities & CL_COMMAND_BUFFER_CAPABILITY_SIMULTANEOUS_USE_KHR) + != 0; + + // If device supports command queue property + // CL_COMMAND_BUFFER_SIMULTANEOUS_USE_KHR test should be skipped + return device_supports_simultaneous_use; + } +}; + +// CL_INCOMPATIBLE_COMMAND_QUEUE_KHR if the properties of any command-queue in +// queues does not contain the minimum properties specified by +// CL_DEVICE_COMMAND_BUFFER_REQUIRED_QUEUE_PROPERTIES_KHR. +struct CreateCommandBufferQueueWithoutMinProperties + : public BasicCommandBufferTest +{ + using BasicCommandBufferTest::BasicCommandBufferTest; + + cl_int Run() override + { + cl_int error = CL_SUCCESS; + + command_buffer = clCreateCommandBufferKHR(1, &queue, nullptr, &error); + test_failure_error_ret(error, CL_INCOMPATIBLE_COMMAND_QUEUE_KHR, + "clCreateCommandBufferKHR should return " + "CL_INCOMPATIBLE_COMMAND_QUEUE_KHR", + TEST_FAIL); + + return CL_SUCCESS; + } + + bool Skip() override + { + cl_command_queue_properties required_properties; + cl_int error = clGetDeviceInfo( + device, CL_DEVICE_COMMAND_BUFFER_REQUIRED_QUEUE_PROPERTIES_KHR, + sizeof(required_properties), &required_properties, NULL); + test_error(error, + "Unable to query " + "CL_DEVICE_COMMAND_BUFFER_REQUIRED_QUEUE_PROPERTIES_KHR"); + + cl_command_queue_properties queue_properties; + error = clGetCommandQueueInfo(queue, CL_QUEUE_PROPERTIES, + sizeof(queue_properties), + &queue_properties, NULL); + test_error(error, "Unable to query CL_QUEUE_PROPERTIES"); + + // Skip if queue properties contains those required + return required_properties == (required_properties & queue_properties); + } +}; + +// CL_INCOMPATIBLE_COMMAND_QUEUE_KHR if any command-queue in queues is an +// out-of-order command-queue and the device associated with the command-queue +// does not support the CL_COMMAND_BUFFER_CAPABILITY_OUT_OF_ORDER_KHR +// capability. +struct CreateCommandBufferDeviceDoesNotSupportOutOfOderQueue + : public BasicCommandBufferTest +{ + CreateCommandBufferDeviceDoesNotSupportOutOfOderQueue( + cl_device_id device, cl_context context, cl_command_queue queue) + : BasicCommandBufferTest(device, context, queue), + out_of_order_queue(nullptr) + {} + + cl_int Run() override + { + cl_int error = CL_SUCCESS; + + command_buffer = + clCreateCommandBufferKHR(1, &out_of_order_queue, nullptr, &error); + test_failure_error_ret(error, CL_INCOMPATIBLE_COMMAND_QUEUE_KHR, + "clCreateCommandBufferKHR should return " + "CL_INCOMPATIBLE_COMMAND_QUEUE_KHR", + TEST_FAIL); + + return CL_SUCCESS; + } + + cl_int SetUp(int elements) override + { + cl_int error = CL_SUCCESS; + + error = BasicCommandBufferTest::SetUp(elements); + test_error(error, "BasicCommandBufferTest::SetUp failed"); + + out_of_order_queue = clCreateCommandQueue( + context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &error); + test_error(error, + "clCreateCommandQueue with " + "CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE failed"); + + return CL_SUCCESS; + } + + bool Skip() override + { + BasicCommandBufferTest::Skip(); + + // If device supports out of order queues test should be skipped + return out_of_order_support != 0; + } + + clCommandQueueWrapper out_of_order_queue; +}; +}; + +int test_negative_create_command_buffer_num_queues(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements) +{ + return MakeAndRunTest(device, context, queue, + num_elements); +} + +int test_negative_create_command_buffer_null_queues(cl_device_id device, + cl_context context, + cl_command_queue queue, + int num_elements) +{ + return MakeAndRunTest(device, context, queue, + num_elements); +} + +int test_negative_create_command_buffer_repeated_properties( + cl_device_id device, cl_context context, cl_command_queue queue, + int num_elements) +{ + return MakeAndRunTest( + device, context, queue, num_elements); +} + +int test_negative_create_command_buffer_not_supported_properties( + cl_device_id device, cl_context context, cl_command_queue queue, + int num_elements) +{ + return MakeAndRunTest( + device, context, queue, num_elements); +} + +int test_negative_create_command_buffer_queue_without_min_properties( + cl_device_id device, cl_context context, cl_command_queue queue, + int num_elements) +{ + return MakeAndRunTest( + device, context, queue, num_elements); +} + +int test_negative_create_command_buffer_device_does_not_support_out_of_order_queue( + cl_device_id device, cl_context context, cl_command_queue queue, + int num_elements) +{ + return MakeAndRunTest< + CreateCommandBufferDeviceDoesNotSupportOutOfOderQueue>( + device, context, queue, num_elements); +} diff --git a/test_conformance/extensions/cl_khr_command_buffer/procs.h b/test_conformance/extensions/cl_khr_command_buffer/procs.h index 12d06d99b..a9829e3df 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/procs.h +++ b/test_conformance/extensions/cl_khr_command_buffer/procs.h @@ -141,6 +141,25 @@ extern int test_finalize_invalid(cl_device_id device, cl_context context, extern int test_finalize_empty(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements); // Command-buffer negative tests +extern int test_negative_create_command_buffer_num_queues( + cl_device_id device, cl_context context, cl_command_queue queue, + int num_elements); +extern int test_negative_create_command_buffer_null_queues( + cl_device_id device, cl_context context, cl_command_queue queue, + int num_elements); +extern int test_negative_create_command_buffer_repeated_properties( + cl_device_id device, cl_context context, cl_command_queue queue, + int num_elements); +extern int test_negative_create_command_buffer_not_supported_properties( + cl_device_id device, cl_context context, cl_command_queue queue, + int num_elements); +extern int test_negative_create_command_buffer_queue_without_min_properties( + cl_device_id device, cl_context context, cl_command_queue queue, + int num_elements); +extern int +test_negative_create_command_buffer_device_does_not_support_out_of_order_queue( + cl_device_id device, cl_context context, cl_command_queue queue, + int num_elements); extern int test_negative_get_command_buffer_info_invalid_command_buffer( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements); @@ -166,4 +185,5 @@ extern int test_negative_get_command_buffer_info_context(cl_device_id device, cl_command_queue queue, int num_elements); + #endif // CL_KHR_COMMAND_BUFFER_PROCS_H From be8b56d94972b4d159e7471aebf5e9ab302d4d81 Mon Sep 17 00:00:00 2001 From: Romaric Jodin <89833130+rjodinchr@users.noreply.github.com> Date: Tue, 16 Apr 2024 17:48:05 +0200 Subject: [PATCH 3/7] Test IMAGE1D_BUFFER in more scenario (#1806) * cl_copy_images * cl_get_info * cl_fill_image * cl_read_write_image * kernel_image_methods * IMAGE1D_BUFFER cannot be created with (USE_|ALLOC_|COPY_)_HOST_PTR * do not allow mipmap with 1D buffer * adjust M to be within maximum_sizes and max_pixels * remove unused variables * make sure M will never be 0 * fix region[0] after refactoring removing mipmap * fix formatting * format with clang-format-11 * fix image1d_buffer creation with gEnablePitch * add missing case in switch * use align_malloc when CL version is at least 2.0 * use CL_DEVICE_NUMERIC_VERSION and align_free * fix free of pitch buffer * fix formatting * fix formatting * fix data->is_aligned --- test_common/harness/imageHelpers.cpp | 15 +- .../images/clCopyImage/CMakeLists.txt | 1 + test_conformance/images/clCopyImage/main.cpp | 31 +- .../clCopyImage/test_copy_1D_buffer.cpp | 516 ++++++++++++++++++ .../images/clCopyImage/test_copy_generic.cpp | 106 +++- .../images/clCopyImage/test_loops.cpp | 34 ++ .../images/clFillImage/CMakeLists.txt | 1 + test_conformance/images/clFillImage/main.cpp | 12 +- .../clFillImage/test_fill_1D_buffer.cpp | 203 +++++++ .../images/clFillImage/test_fill_generic.cpp | 207 +++++-- .../images/clFillImage/test_loops.cpp | 10 + .../images/clGetInfo/CMakeLists.txt | 1 + test_conformance/images/clGetInfo/main.cpp | 12 +- .../images/clGetInfo/test_1D_buffer.cpp | 128 +++++ test_conformance/images/clGetInfo/test_2D.cpp | 31 +- .../images/clGetInfo/test_loops.cpp | 8 + .../images/clReadWriteImage/CMakeLists.txt | 1 + .../images/clReadWriteImage/main.cpp | 12 +- .../images/clReadWriteImage/test_loops.cpp | 9 + .../clReadWriteImage/test_read_1D_buffer.cpp | 256 +++++++++ .../kernel_image_methods/CMakeLists.txt | 1 + .../images/kernel_image_methods/main.cpp | 12 +- .../kernel_image_methods/test_1D_buffer.cpp | 282 ++++++++++ .../kernel_image_methods/test_loops.cpp | 11 + test_conformance/images/testBase.h | 25 +- 25 files changed, 1838 insertions(+), 87 deletions(-) create mode 100644 test_conformance/images/clCopyImage/test_copy_1D_buffer.cpp create mode 100644 test_conformance/images/clFillImage/test_fill_1D_buffer.cpp create mode 100644 test_conformance/images/clGetInfo/test_1D_buffer.cpp create mode 100644 test_conformance/images/clReadWriteImage/test_read_1D_buffer.cpp create mode 100644 test_conformance/images/kernel_image_methods/test_1D_buffer.cpp diff --git a/test_common/harness/imageHelpers.cpp b/test_common/harness/imageHelpers.cpp index 2e0ad742e..52a642383 100644 --- a/test_common/harness/imageHelpers.cpp +++ b/test_common/harness/imageHelpers.cpp @@ -769,10 +769,14 @@ void get_max_sizes( (*numberOfSizes) = 0; - if (image_type == CL_MEM_OBJECT_IMAGE1D) + if (image_type == CL_MEM_OBJECT_IMAGE1D + || image_type == CL_MEM_OBJECT_IMAGE1D_BUFFER) { size_t M = maximum_sizes[0]; + size_t A = max_pixels; + + M = static_cast(fmax(1, fmin(A / M, M))); // Store the size sizes[(*numberOfSizes)][0] = M; @@ -860,6 +864,7 @@ void get_max_sizes( { switch (image_type) { + case CL_MEM_OBJECT_IMAGE1D_BUFFER: case CL_MEM_OBJECT_IMAGE1D: log_info(" size[%d] = [%zu] (%g MB image)\n", j, sizes[j][0], raw_pixel_size * sizes[j][0] * sizes[j][1] @@ -1080,6 +1085,7 @@ cl_ulong get_image_size(image_descriptor const *imageInfo) { switch (imageInfo->type) { + case CL_MEM_OBJECT_IMAGE1D_BUFFER: case CL_MEM_OBJECT_IMAGE1D: imageSize = imageInfo->rowPitch; break; case CL_MEM_OBJECT_IMAGE2D: imageSize = imageInfo->height * imageInfo->rowPitch; @@ -2317,6 +2323,7 @@ int debug_find_vector_in_image(void *imagePtr, image_descriptor *imageInfo, switch (imageInfo->type) { + case CL_MEM_OBJECT_IMAGE1D_BUFFER: case CL_MEM_OBJECT_IMAGE1D: width = (imageInfo->width >> lod) ? (imageInfo->width >> lod) : 1; height = 1; @@ -3513,6 +3520,7 @@ void copy_image_data(image_descriptor *srcImageInfo, switch (srcImageInfo->type) { + case CL_MEM_OBJECT_IMAGE1D_BUFFER: case CL_MEM_OBJECT_IMAGE1D: src_lod = sourcePos[1]; sourcePos_lod[1] = sourcePos_lod[2] = 0; @@ -3558,6 +3566,7 @@ void copy_image_data(image_descriptor *srcImageInfo, size_t dst_height_lod = 1 /*dstImageInfo->height*/; switch (dstImageInfo->type) { + case CL_MEM_OBJECT_IMAGE1D_BUFFER: case CL_MEM_OBJECT_IMAGE1D: dst_lod = destPos[1]; destPos_lod[1] = destPos_lod[2] = 0; @@ -4022,6 +4031,7 @@ cl_ulong compute_mipmapped_image_size(image_descriptor imageInfo) retSize += (cl_ulong)curr_width * curr_height * get_pixel_size(imageInfo.format); break; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: case CL_MEM_OBJECT_IMAGE1D: retSize += (cl_ulong)curr_width * get_pixel_size(imageInfo.format); @@ -4043,6 +4053,7 @@ cl_ulong compute_mipmapped_image_size(image_descriptor imageInfo) case CL_MEM_OBJECT_IMAGE2D: case CL_MEM_OBJECT_IMAGE2D_ARRAY: curr_height = curr_height >> 1 ? curr_height >> 1 : 1; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: case CL_MEM_OBJECT_IMAGE1D: case CL_MEM_OBJECT_IMAGE1D_ARRAY: curr_width = curr_width >> 1 ? curr_width >> 1 : 1; @@ -4080,6 +4091,7 @@ size_t compute_mip_level_offset(image_descriptor *imageInfo, size_t lod) retOffset += (size_t)width * height * get_pixel_size(imageInfo->format); break; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: case CL_MEM_OBJECT_IMAGE1D: retOffset += (size_t)width * get_pixel_size(imageInfo->format); break; @@ -4092,6 +4104,7 @@ size_t compute_mip_level_offset(image_descriptor *imageInfo, size_t lod) case CL_MEM_OBJECT_IMAGE2D: case CL_MEM_OBJECT_IMAGE2D_ARRAY: height = (height >> 1) ? (height >> 1) : 1; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: case CL_MEM_OBJECT_IMAGE1D_ARRAY: case CL_MEM_OBJECT_IMAGE1D: width = (width >> 1) ? (width >> 1) : 1; } diff --git a/test_conformance/images/clCopyImage/CMakeLists.txt b/test_conformance/images/clCopyImage/CMakeLists.txt index bf06dc688..c50000ff7 100644 --- a/test_conformance/images/clCopyImage/CMakeLists.txt +++ b/test_conformance/images/clCopyImage/CMakeLists.txt @@ -10,6 +10,7 @@ set(${MODULE_NAME}_SOURCES test_copy_2D_array.cpp test_copy_3D.cpp test_copy_3D_2D_array.cpp + test_copy_1D_buffer.cpp test_copy_generic.cpp test_loops.cpp ../common.cpp diff --git a/test_conformance/images/clCopyImage/main.cpp b/test_conformance/images/clCopyImage/main.cpp index c2cad0107..5526c3980 100644 --- a/test_conformance/images/clCopyImage/main.cpp +++ b/test_conformance/images/clCopyImage/main.cpp @@ -45,6 +45,21 @@ int test_3D(cl_device_id device, cl_context context, cl_command_queue queue, int { return test_image_set( device, context, queue, k3D ); } +int test_1Dbuffer(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return test_image_set(device, context, queue, k1DBuffer); +} +int test_1DTo1Dbuffer(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return test_image_set(device, context, queue, k1DTo1DBuffer); +} +int test_1DbufferTo1D(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return test_image_set(device, context, queue, k1DBufferTo1D); +} int test_1Darray(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) { return test_image_set( device, context, queue, k1DArray ); @@ -79,17 +94,11 @@ int test_3Dto2Darray(cl_device_id device, cl_context context, cl_command_queue q } test_definition test_list[] = { - ADD_TEST( 1D ), - ADD_TEST( 2D ), - ADD_TEST( 3D ), - ADD_TEST( 1Darray ), - ADD_TEST( 2Darray ), - ADD_TEST( 2Dto3D ), - ADD_TEST( 3Dto2D ), - ADD_TEST( 2Darrayto2D ), - ADD_TEST( 2Dto2Darray ), - ADD_TEST( 2Darrayto3D ), - ADD_TEST( 3Dto2Darray ), + ADD_TEST(1D), ADD_TEST(2D), ADD_TEST(3D), + ADD_TEST(1Darray), ADD_TEST(2Darray), ADD_TEST(2Dto3D), + ADD_TEST(3Dto2D), ADD_TEST(2Darrayto2D), ADD_TEST(2Dto2Darray), + ADD_TEST(2Darrayto3D), ADD_TEST(3Dto2Darray), ADD_TEST(1Dbuffer), + ADD_TEST(1DTo1Dbuffer), ADD_TEST(1DbufferTo1D), }; const int test_num = ARRAY_SIZE( test_list ); diff --git a/test_conformance/images/clCopyImage/test_copy_1D_buffer.cpp b/test_conformance/images/clCopyImage/test_copy_1D_buffer.cpp new file mode 100644 index 000000000..f8f14bd44 --- /dev/null +++ b/test_conformance/images/clCopyImage/test_copy_1D_buffer.cpp @@ -0,0 +1,516 @@ +// +// 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. +// +#include "../testBase.h" + +extern int test_copy_image_generic(cl_context context, cl_command_queue queue, + image_descriptor *srcImageInfo, + image_descriptor *dstImageInfo, + const size_t sourcePos[], + const size_t destPos[], + const size_t regionSize[], MTdata d); + +int test_copy_image_size_1D_buffer(cl_context context, cl_command_queue queue, + image_descriptor *srcImageInfo, + image_descriptor *dstImageInfo, MTdata d) +{ + size_t sourcePos[3], destPos[3], regionSize[3]; + int ret = 0, retCode; + size_t width_lod = srcImageInfo->width; + + // First, try just a full covering region + sourcePos[0] = sourcePos[1] = sourcePos[2] = 0; + destPos[0] = destPos[1] = destPos[2] = 0; + regionSize[0] = srcImageInfo->width; + regionSize[1] = 1; + regionSize[2] = 1; + + retCode = + test_copy_image_generic(context, queue, srcImageInfo, dstImageInfo, + sourcePos, destPos, regionSize, d); + if (retCode < 0) + return retCode; + else + ret += retCode; + + // Now try a sampling of different random regions + for (int i = 0; i < 8; i++) + { + // Pick a random size + regionSize[0] = (width_lod > 8) + ? (size_t)random_in_range(8, (int)width_lod - 1, d) + : width_lod; + + // Now pick positions within valid ranges + sourcePos[0] = (width_lod > regionSize[0]) ? (size_t)random_in_range( + 0, (int)(width_lod - regionSize[0] - 1), d) + : 0; + destPos[0] = (width_lod > regionSize[0]) ? (size_t)random_in_range( + 0, (int)(width_lod - regionSize[0] - 1), d) + : 0; + + + // Go for it! + retCode = + test_copy_image_generic(context, queue, srcImageInfo, srcImageInfo, + sourcePos, destPos, regionSize, d); + if (retCode < 0) + return retCode; + else + ret += retCode; + } + + return ret; +} + +int test_copy_image_set_1D_buffer(cl_device_id device, cl_context context, + cl_command_queue queue, + cl_image_format *format) +{ + size_t maxWidth; + cl_ulong maxAllocSize, memSize; + image_descriptor imageInfo = { 0 }; + RandomSeed seed(gRandomSeed); + size_t pixelSize; + + if (gTestMipmaps) + { + // 1D image buffers don't support mipmaps + // https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_Ext.html#cl_khr_mipmap_image + return 0; + } + + imageInfo.format = format; + imageInfo.height = imageInfo.depth = imageInfo.arraySize = + imageInfo.slicePitch = 0; + imageInfo.type = CL_MEM_OBJECT_IMAGE1D_BUFFER; + pixelSize = get_pixel_size(imageInfo.format); + + int error = clGetDeviceInfo(device, CL_DEVICE_IMAGE_MAX_BUFFER_SIZE, + sizeof(maxWidth), &maxWidth, NULL); + error |= clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, + sizeof(maxAllocSize), &maxAllocSize, NULL); + error |= clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(memSize), + &memSize, NULL); + test_error(error, "Unable to get max image 1D buffer size from device"); + + if (memSize > (cl_ulong)SIZE_MAX) + { + memSize = (cl_ulong)SIZE_MAX; + maxAllocSize = (cl_ulong)SIZE_MAX; + } + + if (gTestSmallImages) + { + for (imageInfo.width = 1; imageInfo.width < 13; imageInfo.width++) + { + size_t rowPadding = gEnablePitch ? 48 : 0; + imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding; + + if (gEnablePitch) + { + do + { + rowPadding++; + imageInfo.rowPitch = + imageInfo.width * pixelSize + rowPadding; + } while ((imageInfo.rowPitch % pixelSize) != 0); + } + + if (gDebugTrace) log_info(" at size %d\n", (int)imageInfo.width); + + int ret = test_copy_image_size_1D_buffer(context, queue, &imageInfo, + &imageInfo, seed); + if (ret) return -1; + } + } + else if (gTestMaxImages) + { + // Try a specific set of maximum sizes + size_t numbeOfSizes; + size_t sizes[100][3]; + + get_max_sizes(&numbeOfSizes, 100, sizes, maxWidth, 1, 1, 1, + maxAllocSize, memSize, CL_MEM_OBJECT_IMAGE1D_BUFFER, + imageInfo.format); + + for (size_t idx = 0; idx < numbeOfSizes; idx++) + { + size_t rowPadding = gEnablePitch ? 48 : 0; + imageInfo.width = sizes[idx][0]; + imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding; + + if (gEnablePitch) + { + do + { + rowPadding++; + imageInfo.rowPitch = + imageInfo.width * pixelSize + rowPadding; + } while ((imageInfo.rowPitch % pixelSize) != 0); + } + + log_info("Testing %d\n", (int)sizes[idx][0]); + if (gDebugTrace) + log_info(" at max size %d\n", (int)sizes[idx][0]); + if (test_copy_image_size_1D_buffer(context, queue, &imageInfo, + &imageInfo, seed)) + return -1; + } + } + else + { + for (int i = 0; i < NUM_IMAGE_ITERATIONS; i++) + { + cl_ulong size; + size_t rowPadding = gEnablePitch ? 48 : 0; + // Loop until we get a size that a) will fit in the max alloc size + // and b) that an allocation of that image, the result array, plus + // offset arrays, will fit in the global ram space + do + { + imageInfo.width = + (size_t)random_log_in_range(16, (int)maxWidth / 32, seed); + + imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding; + + if (gEnablePitch) + { + do + { + rowPadding++; + imageInfo.rowPitch = + imageInfo.width * pixelSize + rowPadding; + } while ((imageInfo.rowPitch % pixelSize) != 0); + } + + size = (size_t)imageInfo.rowPitch * 4; + } while (size > maxAllocSize || (size * 3) > memSize); + + if (gDebugTrace) + { + log_info(" at size %d (row pitch %d) out of %d\n", + (int)imageInfo.width, (int)imageInfo.rowPitch, + (int)maxWidth); + } + + int ret = test_copy_image_size_1D_buffer(context, queue, &imageInfo, + &imageInfo, seed); + if (ret) return -1; + } + } + + return 0; +} + +int test_copy_image_set_1D_1D_buffer(cl_device_id device, cl_context context, + cl_command_queue queue, + cl_image_format *format) +{ + size_t maxWidth; + cl_ulong maxAllocSize, memSize; + image_descriptor imageInfo = { 0 }; + RandomSeed seed(gRandomSeed); + size_t pixelSize; + + if (gTestMipmaps) + { + // 1D image buffers don't support mipmaps + // https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_Ext.html#cl_khr_mipmap_image + return 0; + } + + imageInfo.format = format; + imageInfo.height = imageInfo.depth = imageInfo.arraySize = + imageInfo.slicePitch = 0; + imageInfo.type = CL_MEM_OBJECT_IMAGE1D_BUFFER; + pixelSize = get_pixel_size(imageInfo.format); + + int error = clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_WIDTH, + sizeof(maxWidth), &maxWidth, NULL); + error |= clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, + sizeof(maxAllocSize), &maxAllocSize, NULL); + error |= clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(memSize), + &memSize, NULL); + test_error(error, "Unable to get max image 1D buffer size from device"); + + if (memSize > (cl_ulong)SIZE_MAX) + { + memSize = (cl_ulong)SIZE_MAX; + maxAllocSize = (cl_ulong)SIZE_MAX; + } + + if (gTestSmallImages) + { + for (imageInfo.width = 1; imageInfo.width < 13; imageInfo.width++) + { + size_t rowPadding = gEnablePitch ? 48 : 0; + imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding; + + if (gEnablePitch) + { + do + { + rowPadding++; + imageInfo.rowPitch = + imageInfo.width * pixelSize + rowPadding; + } while ((imageInfo.rowPitch % pixelSize) != 0); + } + + if (gDebugTrace) log_info(" at size %d\n", (int)imageInfo.width); + + image_descriptor srcImageInfo = imageInfo; + srcImageInfo.type = CL_MEM_OBJECT_IMAGE1D; + + int ret = test_copy_image_size_1D_buffer( + context, queue, &srcImageInfo, &imageInfo, seed); + if (ret) return -1; + } + } + else if (gTestMaxImages) + { + // Try a specific set of maximum sizes + size_t numbeOfSizes; + size_t sizes[100][3]; + + get_max_sizes(&numbeOfSizes, 100, sizes, maxWidth, 1, 1, 1, + maxAllocSize, memSize, CL_MEM_OBJECT_IMAGE1D_BUFFER, + imageInfo.format); + + for (size_t idx = 0; idx < numbeOfSizes; idx++) + { + size_t rowPadding = gEnablePitch ? 48 : 0; + imageInfo.width = sizes[idx][0]; + imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding; + + if (gEnablePitch) + { + do + { + rowPadding++; + imageInfo.rowPitch = + imageInfo.width * pixelSize + rowPadding; + } while ((imageInfo.rowPitch % pixelSize) != 0); + } + + log_info("Testing %d\n", (int)sizes[idx][0]); + if (gDebugTrace) + log_info(" at max size %d\n", (int)sizes[idx][0]); + + image_descriptor srcImageInfo = imageInfo; + srcImageInfo.type = CL_MEM_OBJECT_IMAGE1D; + + if (test_copy_image_size_1D_buffer(context, queue, &srcImageInfo, + &imageInfo, seed)) + return -1; + } + } + else + { + for (int i = 0; i < NUM_IMAGE_ITERATIONS; i++) + { + cl_ulong size; + size_t rowPadding = gEnablePitch ? 48 : 0; + // Loop until we get a size that a) will fit in the max alloc size + // and b) that an allocation of that image, the result array, plus + // offset arrays, will fit in the global ram space + do + { + imageInfo.width = + (size_t)random_log_in_range(16, (int)maxWidth / 32, seed); + + imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding; + + if (gEnablePitch) + { + do + { + rowPadding++; + imageInfo.rowPitch = + imageInfo.width * pixelSize + rowPadding; + } while ((imageInfo.rowPitch % pixelSize) != 0); + } + + size = (size_t)imageInfo.rowPitch * 4; + } while (size > maxAllocSize || (size * 3) > memSize); + + if (gDebugTrace) + { + log_info(" at size %d (row pitch %d) out of %d\n", + (int)imageInfo.width, (int)imageInfo.rowPitch, + (int)maxWidth); + } + + image_descriptor srcImageInfo = imageInfo; + srcImageInfo.type = CL_MEM_OBJECT_IMAGE1D; + + int ret = test_copy_image_size_1D_buffer( + context, queue, &srcImageInfo, &imageInfo, seed); + if (ret) return -1; + } + } + + return 0; +} + +int test_copy_image_set_1D_buffer_1D(cl_device_id device, cl_context context, + cl_command_queue queue, + cl_image_format *format) +{ + size_t maxWidth; + cl_ulong maxAllocSize, memSize; + image_descriptor imageInfo = { 0 }; + RandomSeed seed(gRandomSeed); + size_t pixelSize; + + if (gTestMipmaps) + { + // 1D image buffers don't support mipmaps + // https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_Ext.html#cl_khr_mipmap_image + return 0; + } + + imageInfo.format = format; + imageInfo.height = imageInfo.depth = imageInfo.arraySize = + imageInfo.slicePitch = 0; + imageInfo.type = CL_MEM_OBJECT_IMAGE1D_BUFFER; + pixelSize = get_pixel_size(imageInfo.format); + + int error = clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_WIDTH, + sizeof(maxWidth), &maxWidth, NULL); + error |= clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, + sizeof(maxAllocSize), &maxAllocSize, NULL); + error |= clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(memSize), + &memSize, NULL); + test_error(error, "Unable to get max image 1D buffer size from device"); + + if (memSize > (cl_ulong)SIZE_MAX) + { + memSize = (cl_ulong)SIZE_MAX; + maxAllocSize = (cl_ulong)SIZE_MAX; + } + + if (gTestSmallImages) + { + for (imageInfo.width = 1; imageInfo.width < 13; imageInfo.width++) + { + size_t rowPadding = gEnablePitch ? 48 : 0; + imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding; + + if (gEnablePitch) + { + do + { + rowPadding++; + imageInfo.rowPitch = + imageInfo.width * pixelSize + rowPadding; + } while ((imageInfo.rowPitch % pixelSize) != 0); + } + + if (gDebugTrace) log_info(" at size %d\n", (int)imageInfo.width); + + image_descriptor dstImageInfo = imageInfo; + dstImageInfo.type = CL_MEM_OBJECT_IMAGE1D; + + int ret = test_copy_image_size_1D_buffer(context, queue, &imageInfo, + &dstImageInfo, seed); + if (ret) return -1; + } + } + else if (gTestMaxImages) + { + // Try a specific set of maximum sizes + size_t numbeOfSizes; + size_t sizes[100][3]; + + get_max_sizes(&numbeOfSizes, 100, sizes, maxWidth, 1, 1, 1, + maxAllocSize, memSize, CL_MEM_OBJECT_IMAGE1D_BUFFER, + imageInfo.format); + + for (size_t idx = 0; idx < numbeOfSizes; idx++) + { + size_t rowPadding = gEnablePitch ? 48 : 0; + imageInfo.width = sizes[idx][0]; + imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding; + + if (gEnablePitch) + { + do + { + rowPadding++; + imageInfo.rowPitch = + imageInfo.width * pixelSize + rowPadding; + } while ((imageInfo.rowPitch % pixelSize) != 0); + } + + log_info("Testing %d\n", (int)sizes[idx][0]); + if (gDebugTrace) + log_info(" at max size %d\n", (int)sizes[idx][0]); + + image_descriptor dstImageInfo = imageInfo; + dstImageInfo.type = CL_MEM_OBJECT_IMAGE1D; + + if (test_copy_image_size_1D_buffer(context, queue, &imageInfo, + &dstImageInfo, seed)) + return -1; + } + } + else + { + for (int i = 0; i < NUM_IMAGE_ITERATIONS; i++) + { + cl_ulong size; + size_t rowPadding = gEnablePitch ? 48 : 0; + // Loop until we get a size that a) will fit in the max alloc size + // and b) that an allocation of that image, the result array, plus + // offset arrays, will fit in the global ram space + do + { + imageInfo.width = + (size_t)random_log_in_range(16, (int)maxWidth / 32, seed); + + imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding; + + if (gEnablePitch) + { + do + { + rowPadding++; + imageInfo.rowPitch = + imageInfo.width * pixelSize + rowPadding; + } while ((imageInfo.rowPitch % pixelSize) != 0); + } + + size = (size_t)imageInfo.rowPitch * 4; + } while (size > maxAllocSize || (size * 3) > memSize); + + if (gDebugTrace) + { + log_info(" at size %d (row pitch %d) out of %d\n", + (int)imageInfo.width, (int)imageInfo.rowPitch, + (int)maxWidth); + } + + image_descriptor dstImageInfo = imageInfo; + dstImageInfo.type = CL_MEM_OBJECT_IMAGE1D; + + int ret = test_copy_image_size_1D_buffer(context, queue, &imageInfo, + &dstImageInfo, seed); + if (ret) return -1; + } + } + + return 0; +} diff --git a/test_conformance/images/clCopyImage/test_copy_generic.cpp b/test_conformance/images/clCopyImage/test_copy_generic.cpp index 888ca6ec5..54a156fed 100644 --- a/test_conformance/images/clCopyImage/test_copy_generic.cpp +++ b/test_conformance/images/clCopyImage/test_copy_generic.cpp @@ -14,12 +14,18 @@ // limitations under the License. // #include "../testBase.h" +#include static void CL_CALLBACK free_pitch_buffer( cl_mem image, void *buf ) { free( buf ); } +static void CL_CALLBACK release_cl_buffer(cl_mem image, void *buf) +{ + clReleaseMemObject((cl_mem)buf); +} + cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr& data, image_descriptor *imageInfo, int *error ) { cl_mem img; @@ -69,6 +75,71 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr if ( gEnablePitch ) host_ptr = malloc( imageInfo->arraySize * imageInfo->slicePitch ); break; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: + if (gDebugTrace) + log_info(" - Creating 1D buffer image %d ...\n", + (int)imageInfo->width); + { + cl_int err; + cl_mem_flags buffer_flags = CL_MEM_READ_WRITE; + if (gEnablePitch) + { + cl_device_id device; + err = + clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE, + sizeof(device), &device, nullptr); + if (err != CL_SUCCESS) + { + log_error( + "Error: Could not get CL_QUEUE_DEVICE from queue"); + return NULL; + } + char major_version; + err = clGetDeviceInfo(device, CL_DEVICE_VERSION, + sizeof(major_version), &major_version, + nullptr); + if (err != CL_SUCCESS) + { + log_error("Error: Could not get CL_DEVICE_VERSION from " + "device"); + return NULL; + } + if (major_version == '1') + { + host_ptr = malloc(imageInfo->rowPitch); + } + else + { + cl_uint base_address_alignment = 0; + err = clGetDeviceInfo( + device, CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT, + sizeof(base_address_alignment), + &base_address_alignment, nullptr); + if (err != CL_SUCCESS) + { + log_error("ERROR: Could not get " + "CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT " + "from device"); + return NULL; + } + host_ptr = align_malloc(imageInfo->rowPitch, + base_address_alignment); + } + buffer_flags |= CL_MEM_USE_HOST_PTR; + } + + cl_mem buffer = clCreateBuffer(context, buffer_flags, + imageInfo->rowPitch, NULL, &err); + if (err != CL_SUCCESS) + { + log_error("ERROR: Could not create buffer for 1D buffer " + "image. %ld bytes\n", + imageInfo->width); + return NULL; + } + imageDesc.buffer = buffer; + } + break; } if ( gDebugTrace && gTestMipmaps ) @@ -81,7 +152,10 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr log_error( "ERROR: Unable to create backing store for pitched 3D image. %ld bytes\n", imageInfo->depth * imageInfo->slicePitch ); return NULL; } - mem_flags = CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR; + if (imageInfo->type != CL_MEM_OBJECT_IMAGE1D_BUFFER) + { + mem_flags = CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR; + } } img = clCreateImage(context, mem_flags, imageInfo->format, &imageDesc, host_ptr, error); @@ -103,6 +177,21 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr free(host_ptr); } + if (imageDesc.buffer != NULL) + { + int callbackError = clSetMemObjectDestructorCallback( + img, release_cl_buffer, imageDesc.buffer); + if (callbackError != CL_SUCCESS) + { + log_error("Error: Unable to attach destructor callback to 1d " + "buffer image. Err: %d\n", + callbackError); + clReleaseMemObject(imageDesc.buffer); + clReleaseMemObject(img); + return NULL; + } + } + if ( *error != CL_SUCCESS ) { long long unsigned imageSize = get_image_size_mb(imageInfo); @@ -141,6 +230,12 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr (int)imageInfo->arraySize, imageSize, IGetErrorString(*error)); break; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: + log_error( + "ERROR: Unable to create 1D buffer image of size %d (%llu " + "MB):(%s)", + (int)imageInfo->width, imageSize, IGetErrorString(*error)); + break; } log_error("ERROR: and %llu mip levels\n", (unsigned long long) imageInfo->num_mip_levels); return NULL; @@ -161,6 +256,7 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr height = imageInfo->arraySize; depth = 1; break; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: case CL_MEM_OBJECT_IMAGE1D: height = depth = 1; break; @@ -196,6 +292,7 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr case CL_MEM_OBJECT_IMAGE1D_ARRAY: origin[ 2 ] = lod; break; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: case CL_MEM_OBJECT_IMAGE1D: origin[ 1 ] = lod; break; @@ -210,6 +307,7 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr case CL_MEM_OBJECT_IMAGE2D: height = ( imageInfo->height >> lod ) ? (imageInfo->height >> lod) : 1; case CL_MEM_OBJECT_IMAGE1D_ARRAY: + case CL_MEM_OBJECT_IMAGE1D_BUFFER: case CL_MEM_OBJECT_IMAGE1D: width = ( imageInfo->width >> lod ) ? (imageInfo->width >> lod) : 1; } @@ -260,6 +358,7 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr break; case CL_MEM_OBJECT_IMAGE1D_ARRAY: case CL_MEM_OBJECT_IMAGE1D: + case CL_MEM_OBJECT_IMAGE1D_BUFFER: dstPitch2D = mappedSlice; break; } @@ -391,8 +490,7 @@ int test_copy_image_generic( cl_context context, cl_command_queue queue, image_d switch(dstImageInfo->type) { case CL_MEM_OBJECT_IMAGE1D: - dst_lod = destPos[1]; - break; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: dst_lod = destPos[1]; break; case CL_MEM_OBJECT_IMAGE1D_ARRAY: case CL_MEM_OBJECT_IMAGE2D: dst_lod = destPos[2]; @@ -407,6 +505,7 @@ int test_copy_image_generic( cl_context context, cl_command_queue queue, image_d } switch (dstImageInfo->type) { + case CL_MEM_OBJECT_IMAGE1D_BUFFER: case CL_MEM_OBJECT_IMAGE1D: if( gTestMipmaps ) origin[ 1 ] = dst_lod; @@ -542,6 +641,7 @@ int test_copy_image_generic( cl_context context, cl_command_queue queue, image_d secondDim = dstImageInfo->height; break; } + case CL_MEM_OBJECT_IMAGE1D_BUFFER: case CL_MEM_OBJECT_IMAGE1D: { break; } diff --git a/test_conformance/images/clCopyImage/test_loops.cpp b/test_conformance/images/clCopyImage/test_loops.cpp index e839cfdf3..ea60d3560 100644 --- a/test_conformance/images/clCopyImage/test_loops.cpp +++ b/test_conformance/images/clCopyImage/test_loops.cpp @@ -24,6 +24,18 @@ extern int test_copy_image_set_2D_array( cl_device_id device, cl_context context extern int test_copy_image_set_2D_3D( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, bool reverse ); extern int test_copy_image_set_2D_2D_array( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, bool reverse ); extern int test_copy_image_set_3D_2D_array( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, bool reverse ); +extern int test_copy_image_set_1D_buffer(cl_device_id device, + cl_context context, + cl_command_queue queue, + cl_image_format *format); +extern int test_copy_image_set_1D_1D_buffer(cl_device_id device, + cl_context context, + cl_command_queue queue, + cl_image_format *format); +extern int test_copy_image_set_1D_buffer_1D(cl_device_id device, + cl_context context, + cl_command_queue queue, + cl_image_format *format); int test_image_type( cl_device_id device, cl_context context, cl_command_queue queue, MethodsToTest testMethod, cl_mem_flags flags ) { @@ -87,6 +99,18 @@ int test_image_type( cl_device_id device, cl_context context, cl_command_queue q name = "3D -> 2D array"; imageType = CL_MEM_OBJECT_IMAGE3D; break; + case k1DBuffer: + name = "1D buffer -> 1D buffer"; + imageType = CL_MEM_OBJECT_IMAGE1D_BUFFER; + break; + case k1DTo1DBuffer: + name = "1D -> 1D buffer"; + imageType = CL_MEM_OBJECT_IMAGE1D_BUFFER; + break; + case k1DBufferTo1D: + name = "1D buffer -> 1D"; + imageType = CL_MEM_OBJECT_IMAGE1D_BUFFER; + break; } if(gTestMipmaps) @@ -138,6 +162,16 @@ int test_image_type( cl_device_id device, cl_context context, cl_command_queue q test_return = test_copy_image_set_3D_2D_array( device, context, queue, &formatList[ i ], true); else if( testMethod == k3DTo2DArray) test_return = test_copy_image_set_3D_2D_array( device, context, queue, &formatList[ i ], false); + else if (testMethod == k1DBuffer) + test_return = test_copy_image_set_1D_buffer(device, context, queue, + &formatList[i]); + else if (testMethod == k1DBufferTo1D) + test_return = test_copy_image_set_1D_buffer_1D( + device, context, queue, &formatList[i]); + else if (testMethod == k1DTo1DBuffer) + test_return = test_copy_image_set_1D_1D_buffer( + device, context, queue, &formatList[i]); + if (test_return) { gFailCount++; diff --git a/test_conformance/images/clFillImage/CMakeLists.txt b/test_conformance/images/clFillImage/CMakeLists.txt index a4de32128..974aff291 100644 --- a/test_conformance/images/clFillImage/CMakeLists.txt +++ b/test_conformance/images/clFillImage/CMakeLists.txt @@ -4,6 +4,7 @@ set(MODULE_NAME CL_FILL_IMAGES) set(${MODULE_NAME}_SOURCES main.cpp test_fill_1D.cpp + test_fill_1D_buffer.cpp test_fill_1D_array.cpp test_fill_2D.cpp test_fill_2D_array.cpp diff --git a/test_conformance/images/clFillImage/main.cpp b/test_conformance/images/clFillImage/main.cpp index b19d85afd..62e62c63c 100644 --- a/test_conformance/images/clFillImage/main.cpp +++ b/test_conformance/images/clFillImage/main.cpp @@ -52,13 +52,15 @@ int test_2Darray(cl_device_id device, cl_context context, cl_command_queue queue { return test_image_set(device, context, queue, k2DArray); } +int test_1Dbuffer(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return test_image_set(device, context, queue, k1DBuffer); +} test_definition test_list[] = { - ADD_TEST( 1D ), - ADD_TEST( 2D ), - ADD_TEST( 3D ), - ADD_TEST( 1Darray ), - ADD_TEST( 2Darray ), + ADD_TEST(1D), ADD_TEST(2D), ADD_TEST(3D), + ADD_TEST(1Darray), ADD_TEST(2Darray), ADD_TEST(1Dbuffer), }; const int test_num = ARRAY_SIZE( test_list ); diff --git a/test_conformance/images/clFillImage/test_fill_1D_buffer.cpp b/test_conformance/images/clFillImage/test_fill_1D_buffer.cpp new file mode 100644 index 000000000..e3079c467 --- /dev/null +++ b/test_conformance/images/clFillImage/test_fill_1D_buffer.cpp @@ -0,0 +1,203 @@ +// +// 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. +// +#include "../testBase.h" + +// Defined in test_fill_2D_3D.cpp +extern int test_fill_image_generic(cl_context context, cl_command_queue queue, + image_descriptor *imageInfo, + const size_t origin[], const size_t region[], + ExplicitType outputType, MTdata d); + + +int test_fill_image_size_1D_buffer(cl_context context, cl_command_queue queue, + image_descriptor *imageInfo, + ExplicitType outputType, MTdata d) +{ + size_t origin[3], region[3]; + int ret = 0, retCode; + + // First, try just a full covering region fill + origin[0] = origin[1] = origin[2] = 0; + region[0] = imageInfo->width; + region[1] = 1; + region[2] = 1; + + retCode = test_fill_image_generic(context, queue, imageInfo, origin, region, + outputType, d); + if (retCode < 0) + return retCode; + else + ret += retCode; + + // Now try a sampling of different random regions + for (int i = 0; i < 8; i++) + { + // Pick a random size + region[0] = (imageInfo->width > 8) + ? (size_t)random_in_range(8, (int)imageInfo->width - 1, d) + : imageInfo->width; + + // Now pick positions within valid ranges + origin[0] = (imageInfo->width > region[0]) ? (size_t)random_in_range( + 0, (int)(imageInfo->width - region[0] - 1), d) + : 0; + + // Go for it! + retCode = test_fill_image_generic(context, queue, imageInfo, origin, + region, outputType, d); + if (retCode < 0) + return retCode; + else + ret += retCode; + } + + return ret; +} + + +int test_fill_image_set_1D_buffer(cl_device_id device, cl_context context, + cl_command_queue queue, + cl_image_format *format, + ExplicitType outputType) +{ + size_t maxWidth; + cl_ulong maxAllocSize, memSize; + image_descriptor imageInfo = { 0 }; + RandomSeed seed(gRandomSeed); + const size_t rowPadding_default = 48; + size_t rowPadding = gEnablePitch ? rowPadding_default : 0; + size_t pixelSize; + + memset(&imageInfo, 0x0, sizeof(image_descriptor)); + imageInfo.type = CL_MEM_OBJECT_IMAGE1D_BUFFER; + imageInfo.format = format; + pixelSize = get_pixel_size(imageInfo.format); + + int error = clGetDeviceInfo(device, CL_DEVICE_IMAGE_MAX_BUFFER_SIZE, + sizeof(maxWidth), &maxWidth, NULL); + error |= clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, + sizeof(maxAllocSize), &maxAllocSize, NULL); + error |= clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(memSize), + &memSize, NULL); + test_error(error, "Unable to get max image 2D size from device"); + + if (memSize > (cl_ulong)SIZE_MAX) + { + memSize = (cl_ulong)SIZE_MAX; + maxAllocSize = (cl_ulong)SIZE_MAX; + } + + if (gTestSmallImages) + { + for (imageInfo.width = 1; imageInfo.width < 13; imageInfo.width++) + { + imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding; + + if (gEnablePitch) + { + rowPadding = rowPadding_default; + do + { + rowPadding++; + imageInfo.rowPitch = + imageInfo.width * pixelSize + rowPadding; + } while ((imageInfo.rowPitch % pixelSize) != 0); + } + + if (gDebugTrace) + log_info(" at size %d,%d\n", (int)imageInfo.width, + (int)imageInfo.height); + + int ret = test_fill_image_size_1D_buffer(context, queue, &imageInfo, + outputType, seed); + if (ret) return -1; + } + } + else if (gTestMaxImages) + { + // Try a specific set of maximum sizes + size_t numbeOfSizes; + size_t sizes[100][3]; + + get_max_sizes(&numbeOfSizes, 100, sizes, maxWidth, 1, 1, 1, + maxAllocSize, memSize, CL_MEM_OBJECT_IMAGE1D_BUFFER, + imageInfo.format); + + for (size_t idx = 0; idx < numbeOfSizes; idx++) + { + imageInfo.width = sizes[idx][0]; + imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding; + + if (gEnablePitch) + { + rowPadding = rowPadding_default; + do + { + rowPadding++; + imageInfo.rowPitch = + imageInfo.width * pixelSize + rowPadding; + } while ((imageInfo.rowPitch % pixelSize) != 0); + } + + log_info("Testing %d\n", (int)sizes[idx][0]); + if (gDebugTrace) + log_info(" at max size %d\n", (int)sizes[idx][0]); + if (test_fill_image_size_1D_buffer(context, queue, &imageInfo, + outputType, seed)) + return -1; + } + } + else + { + for (int i = 0; i < NUM_IMAGE_ITERATIONS; i++) + { + cl_ulong size; + // Loop until we get a size that a) will fit in the max alloc size + // and b) that an allocation of that image, the result array, plus + // offset arrays, will fit in the global ram space + do + { + imageInfo.width = + (size_t)random_log_in_range(16, (int)maxWidth / 32, seed); + + imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding; + + if (gEnablePitch) + { + rowPadding = rowPadding_default; + do + { + rowPadding++; + imageInfo.rowPitch = + imageInfo.width * pixelSize + rowPadding; + } while ((imageInfo.rowPitch % pixelSize) != 0); + } + + size = (size_t)imageInfo.rowPitch * 4; + } while (size > maxAllocSize || (size * 3) > memSize); + + if (gDebugTrace) + log_info(" at size %d (row pitch %d) out of %d\n", + (int)imageInfo.width, (int)imageInfo.rowPitch, + (int)maxWidth); + int ret = test_fill_image_size_1D_buffer(context, queue, &imageInfo, + outputType, seed); + if (ret) return -1; + } + } + + return 0; +} diff --git a/test_conformance/images/clFillImage/test_fill_generic.cpp b/test_conformance/images/clFillImage/test_fill_generic.cpp index 6cd6beb0e..ca9a1bfa3 100644 --- a/test_conformance/images/clFillImage/test_fill_generic.cpp +++ b/test_conformance/images/clFillImage/test_fill_generic.cpp @@ -17,9 +17,27 @@ extern void read_image_pixel_float( void *imageData, image_descriptor *imageInfo, int x, int y, int z, float *outData ); -static void CL_CALLBACK free_pitch_buffer( cl_mem image, void *buf ) +struct pitch_buffer_data { - free( buf ); + void *buf; + bool is_aligned; +}; +static void CL_CALLBACK free_pitch_buffer(cl_mem image, void *data) +{ + struct pitch_buffer_data *d = (struct pitch_buffer_data *)data; + if (d->is_aligned) + { + align_free(d->buf); + } + else + { + free(d->buf); + } + free(d); +} +static void CL_CALLBACK release_cl_buffer(cl_mem image, void *buf) +{ + clReleaseMemObject((cl_mem)buf); } cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr& data, image_descriptor *imageInfo, int *error ) @@ -38,6 +56,26 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr imageDesc.image_row_pitch = gEnablePitch ? imageInfo->rowPitch : 0; imageDesc.image_slice_pitch = gEnablePitch ? imageInfo->slicePitch : 0; + cl_version version; + cl_device_id device; + { + cl_int err = clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE, + sizeof(device), &device, nullptr); + if (err != CL_SUCCESS) + { + log_error("Error: Could not get CL_QUEUE_DEVICE from queue"); + return NULL; + } + err = clGetDeviceInfo(device, CL_DEVICE_NUMERIC_VERSION, + sizeof(version), &version, nullptr); + if (err != CL_SUCCESS) + { + log_error("Error: Could not get CL_DEVICE_NUMERIC_VERSION from " + "device"); + return NULL; + } + } + switch (imageInfo->type) { case CL_MEM_OBJECT_IMAGE1D: @@ -70,6 +108,50 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr if ( gEnablePitch ) host_ptr = malloc( imageInfo->arraySize * imageInfo->slicePitch ); break; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: + if (gDebugTrace) + log_info(" - Creating 1D buffer image %d ...\n", + (int)imageInfo->width); + { + cl_int err; + cl_mem_flags buffer_flags = CL_MEM_READ_WRITE; + if (gEnablePitch) + { + if (CL_VERSION_MAJOR(version) == 1) + { + host_ptr = malloc(imageInfo->rowPitch); + } + else + { + cl_uint base_address_alignment = 0; + err = clGetDeviceInfo( + device, CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT, + sizeof(base_address_alignment), + &base_address_alignment, nullptr); + if (err != CL_SUCCESS) + { + log_error("ERROR: Could not get " + "CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT " + "from device"); + return NULL; + } + host_ptr = align_malloc(imageInfo->rowPitch, + base_address_alignment); + } + buffer_flags |= CL_MEM_USE_HOST_PTR; + } + cl_mem buffer = clCreateBuffer( + context, buffer_flags, imageInfo->rowPitch, host_ptr, &err); + if (err != CL_SUCCESS) + { + log_error("ERROR: Could not create buffer for 1D buffer " + "image. %ld bytes\n", + imageInfo->rowPitch); + return NULL; + } + imageDesc.buffer = buffer; + } + break; } if (gEnablePitch) @@ -79,26 +161,63 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr log_error( "ERROR: Unable to create backing store for pitched 3D image. %ld bytes\n", imageInfo->depth * imageInfo->slicePitch ); return NULL; } - mem_flags = CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR; + if (imageInfo->type != CL_MEM_OBJECT_IMAGE1D_BUFFER) + { + mem_flags = CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR; + } } - img = clCreateImage(context, mem_flags, imageInfo->format, &imageDesc, host_ptr, error); + if (imageInfo->type != CL_MEM_OBJECT_IMAGE1D_BUFFER) + { + img = clCreateImage(context, mem_flags, imageInfo->format, &imageDesc, + host_ptr, error); + } + else + { + img = clCreateImage(context, mem_flags, imageInfo->format, &imageDesc, + nullptr, error); + } if (gEnablePitch) { - if ( *error == CL_SUCCESS ) + struct pitch_buffer_data *data = (struct pitch_buffer_data *)malloc( + sizeof(struct pitch_buffer_data)); + data->buf = host_ptr; + data->is_aligned = (CL_VERSION_MAJOR(version) != 1) + && (imageInfo->type == CL_MEM_OBJECT_IMAGE1D_BUFFER); + if (*error == CL_SUCCESS) { - int callbackError = clSetMemObjectDestructorCallback( img, free_pitch_buffer, host_ptr ); - if ( CL_SUCCESS != callbackError ) + int callbackError = + clSetMemObjectDestructorCallback(img, free_pitch_buffer, data); + if (CL_SUCCESS != callbackError) { - free( host_ptr ); - log_error( "ERROR: Unable to attach destructor callback to pitched 3D image. Err: %d\n", callbackError ); - clReleaseMemObject( img ); + free_pitch_buffer(img, data); + log_error("ERROR: Unable to attach destructor callback to " + "pitched 3D image. Err: %d\n", + callbackError); + clReleaseMemObject(img); return NULL; } } else - free(host_ptr); + { + free_pitch_buffer(img, data); + } + } + + if (imageDesc.buffer != NULL) + { + int callbackError = clSetMemObjectDestructorCallback( + img, release_cl_buffer, imageDesc.buffer); + if (callbackError != CL_SUCCESS) + { + log_error("Error: Unable to attach destructor callback to 1d " + "buffer image. Err: %d\n", + callbackError); + clReleaseMemObject(imageDesc.buffer); + clReleaseMemObject(img); + return NULL; + } } if ( *error != CL_SUCCESS ) @@ -122,6 +241,12 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr case CL_MEM_OBJECT_IMAGE2D_ARRAY: log_error( "ERROR: Unable to create 2D image array of size %d x %d x %d (%llu MB): %s\n", (int)imageInfo->width, (int)imageInfo->height, (int)imageInfo->arraySize, imageSize, IGetErrorString( *error ) ); break; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: + log_error( + "ERROR: Unable to create 1D buffer image of size %d (%llu " + "MB):(%s)", + (int)imageInfo->width, imageSize, IGetErrorString(*error)); + break; } return NULL; } @@ -139,6 +264,7 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr depth = 1; imageSize = imageInfo->rowPitch * imageInfo->arraySize; break; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: case CL_MEM_OBJECT_IMAGE1D: height = depth = 1; imageSize = imageInfo->rowPitch; @@ -194,8 +320,7 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr break; case CL_MEM_OBJECT_IMAGE1D_ARRAY: case CL_MEM_OBJECT_IMAGE1D: - dstPitch2D = mappedSlice; - break; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: dstPitch2D = mappedSlice; break; } for ( size_t z = 0; z < depth; z++ ) @@ -286,6 +411,9 @@ int test_fill_image_generic( cl_context context, cl_command_queue queue, image_d case CL_MEM_OBJECT_IMAGE2D_ARRAY: dataBytes = imageInfo->arraySize * imageInfo->slicePitch; break; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: + dataBytes = imageInfo->rowPitch; + break; } if (dataBytes > imgData.getSize()) @@ -398,6 +526,7 @@ int test_fill_image_generic( cl_context context, cl_command_queue queue, image_d size_t imageRegion[ 3 ] = { imageInfo->width, 1, 1 }; switch (imageInfo->type) { + case CL_MEM_OBJECT_IMAGE1D_BUFFER: case CL_MEM_OBJECT_IMAGE1D: break; case CL_MEM_OBJECT_IMAGE2D: @@ -437,28 +566,30 @@ int test_fill_image_generic( cl_context context, cl_command_queue queue, image_d size_t secondDim = 1; switch (imageInfo->type) { - case CL_MEM_OBJECT_IMAGE1D: - secondDim = 1; - thirdDim = 1; - break; - case CL_MEM_OBJECT_IMAGE2D: - secondDim = imageInfo->height; - thirdDim = 1; - break; - case CL_MEM_OBJECT_IMAGE3D: - secondDim = imageInfo->height; - thirdDim = imageInfo->depth; - break; - case CL_MEM_OBJECT_IMAGE1D_ARRAY: - secondDim = imageInfo->arraySize; - thirdDim = 1; - break; - case CL_MEM_OBJECT_IMAGE2D_ARRAY: - secondDim = imageInfo->height; - thirdDim = imageInfo->arraySize; - break; - default: - log_error("Test error: unhandled image type at %s:%d\n",__FILE__,__LINE__); + case CL_MEM_OBJECT_IMAGE1D_BUFFER: + case CL_MEM_OBJECT_IMAGE1D: + secondDim = 1; + thirdDim = 1; + break; + case CL_MEM_OBJECT_IMAGE2D: + secondDim = imageInfo->height; + thirdDim = 1; + break; + case CL_MEM_OBJECT_IMAGE3D: + secondDim = imageInfo->height; + thirdDim = imageInfo->depth; + break; + case CL_MEM_OBJECT_IMAGE1D_ARRAY: + secondDim = imageInfo->arraySize; + thirdDim = 1; + break; + case CL_MEM_OBJECT_IMAGE2D_ARRAY: + secondDim = imageInfo->height; + thirdDim = imageInfo->arraySize; + break; + default: + log_error("Test error: unhandled image type at %s:%d\n", __FILE__, + __LINE__); }; // Count the number of bytes successfully matched @@ -485,8 +616,10 @@ int test_fill_image_generic( cl_context context, cl_command_queue queue, image_d total_matched += scanlineSize; sourcePtr += imageInfo->rowPitch; - if((imageInfo->type == CL_MEM_OBJECT_IMAGE1D_ARRAY || imageInfo->type == CL_MEM_OBJECT_IMAGE1D)) - destPtr += mappedSlice; + if ((imageInfo->type == CL_MEM_OBJECT_IMAGE1D_ARRAY + || imageInfo->type == CL_MEM_OBJECT_IMAGE1D + || imageInfo->type == CL_MEM_OBJECT_IMAGE1D_BUFFER)) + destPtr += mappedSlice; else destPtr += mappedRow; } diff --git a/test_conformance/images/clFillImage/test_loops.cpp b/test_conformance/images/clFillImage/test_loops.cpp index 126ea0eba..95677aec6 100644 --- a/test_conformance/images/clFillImage/test_loops.cpp +++ b/test_conformance/images/clFillImage/test_loops.cpp @@ -23,6 +23,11 @@ extern int test_fill_image_set_2D( cl_device_id device, cl_context context, cl_c extern int test_fill_image_set_3D( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, ExplicitType outputType ); extern int test_fill_image_set_1D_array( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, ExplicitType outputType ); extern int test_fill_image_set_2D_array( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, ExplicitType outputType ); +extern int test_fill_image_set_1D_buffer(cl_device_id device, + cl_context context, + cl_command_queue queue, + cl_image_format *format, + ExplicitType outputType); typedef int (*test_func)(cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, ExplicitType outputType); @@ -60,6 +65,11 @@ int test_image_type( cl_device_id device, cl_context context, cl_command_queue q imageType = CL_MEM_OBJECT_IMAGE3D; test_fn = &test_fill_image_set_3D; break; + case k1DBuffer: + name = "1D Image Buffer Fill"; + imageType = CL_MEM_OBJECT_IMAGE1D_BUFFER; + test_fn = &test_fill_image_set_1D_buffer; + break; default: log_error("Unhandled method\n"); return -1; } diff --git a/test_conformance/images/clGetInfo/CMakeLists.txt b/test_conformance/images/clGetInfo/CMakeLists.txt index f71c49ceb..c59f69762 100644 --- a/test_conformance/images/clGetInfo/CMakeLists.txt +++ b/test_conformance/images/clGetInfo/CMakeLists.txt @@ -7,6 +7,7 @@ set(${MODULE_NAME}_SOURCES test_2D.cpp test_loops.cpp test_3D.cpp + test_1D_buffer.cpp ../common.cpp ) diff --git a/test_conformance/images/clGetInfo/main.cpp b/test_conformance/images/clGetInfo/main.cpp index 80b3cbb20..933eb00fc 100644 --- a/test_conformance/images/clGetInfo/main.cpp +++ b/test_conformance/images/clGetInfo/main.cpp @@ -54,13 +54,15 @@ int test_2Darray(cl_device_id device, cl_context context, cl_command_queue queue { return test_image_set( device, context, CL_MEM_OBJECT_IMAGE2D_ARRAY ); } +int test_1Dbuffer(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return test_image_set(device, context, CL_MEM_OBJECT_IMAGE1D_BUFFER); +} test_definition test_list[] = { - ADD_TEST( 1D ), - ADD_TEST( 2D ), - ADD_TEST( 3D ), - ADD_TEST( 1Darray ), - ADD_TEST( 2Darray ), + ADD_TEST(1D), ADD_TEST(2D), ADD_TEST(3D), + ADD_TEST(1Darray), ADD_TEST(2Darray), ADD_TEST(1Dbuffer), }; const int test_num = ARRAY_SIZE( test_list ); diff --git a/test_conformance/images/clGetInfo/test_1D_buffer.cpp b/test_conformance/images/clGetInfo/test_1D_buffer.cpp new file mode 100644 index 000000000..7b0c2d1c6 --- /dev/null +++ b/test_conformance/images/clGetInfo/test_1D_buffer.cpp @@ -0,0 +1,128 @@ +// +// 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. +// +#include "../testBase.h" +#include + +extern int test_get_image_info_single(cl_context context, + image_descriptor *imageInfo, MTdata d, + cl_mem_flags flags, size_t row_pitch, + size_t slice_pitch); + + +int test_get_image_info_1D_buffer(cl_device_id device, cl_context context, + cl_image_format *format, cl_mem_flags flags) +{ + size_t maxWidth; + cl_ulong maxAllocSize, memSize; + image_descriptor imageInfo = { 0 }; + RandomSeed seed(gRandomSeed); + size_t pixelSize; + + memset(&imageInfo, 0x0, sizeof(image_descriptor)); + imageInfo.type = CL_MEM_OBJECT_IMAGE1D_BUFFER; + imageInfo.format = format; + pixelSize = get_pixel_size(imageInfo.format); + + int error = clGetDeviceInfo(device, CL_DEVICE_IMAGE_MAX_BUFFER_SIZE, + sizeof(maxWidth), &maxWidth, NULL); + error |= clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, + sizeof(maxAllocSize), &maxAllocSize, NULL); + error |= clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(memSize), + &memSize, NULL); + test_error(error, "Unable to get max image 1D size from device"); + + if (memSize > (cl_ulong)SIZE_MAX) + { + memSize = (cl_ulong)SIZE_MAX; + maxAllocSize = (cl_ulong)SIZE_MAX; + } + + if (gTestSmallImages) + { + for (imageInfo.width = 1; imageInfo.width < 13; imageInfo.width++) + { + imageInfo.rowPitch = imageInfo.width * pixelSize; + if (gDebugTrace) + log_info(" at size %d (flags 0x%x pitch %d)\n", + (int)imageInfo.width, (unsigned int)flags, + (int)imageInfo.rowPitch); + if (test_get_image_info_single(context, &imageInfo, seed, flags, 0, + 0)) + return -1; + } + } + else if (gTestMaxImages) + { + // Try a specific set of maximum sizes + size_t numbeOfSizes; + size_t sizes[100][3]; + + get_max_sizes(&numbeOfSizes, 100, sizes, maxWidth, 1, 1, 1, + maxAllocSize, memSize, CL_MEM_OBJECT_IMAGE1D_BUFFER, + imageInfo.format); + + for (size_t idx = 0; idx < numbeOfSizes; idx++) + { + imageInfo.width = sizes[idx][0]; + imageInfo.rowPitch = imageInfo.width * pixelSize; + log_info("Testing %d x 1\n", (int)sizes[idx][0]); + if (gDebugTrace) + log_info(" at max size %d (flags 0x%x pitch %d)\n", + (int)imageInfo.width, (unsigned int)flags, + (int)imageInfo.rowPitch); + if (test_get_image_info_single(context, &imageInfo, seed, flags, 0, + 0)) + return -1; + } + } + else + { + for (int i = 0; i < NUM_IMAGE_ITERATIONS; i++) + { + cl_ulong size; + // Loop until we get a size that a) will fit in the max alloc size + // and b) that an allocation of that image, the result array, plus + // offset arrays, will fit in the global ram space + do + { + imageInfo.width = + (size_t)random_log_in_range(16, (int)maxWidth / 32, seed); + + imageInfo.rowPitch = imageInfo.width * pixelSize; + size_t extraWidth = (int)random_log_in_range(0, 64, seed); + imageInfo.rowPitch += extraWidth; + + do + { + extraWidth++; + imageInfo.rowPitch += extraWidth; + } while ((imageInfo.rowPitch % pixelSize) != 0); + + size = (cl_ulong)imageInfo.rowPitch * 4; + } while (size > maxAllocSize || (size * 3) > memSize); + + if (gDebugTrace) + log_info(" at size %d (flags 0x%x pitch %d) out of %d\n", + (int)imageInfo.width, (unsigned int)flags, + (int)imageInfo.rowPitch, (int)maxWidth); + if (test_get_image_info_single(context, &imageInfo, seed, flags, 0, + 0)) + return -1; + } + } + + return 0; +} diff --git a/test_conformance/images/clGetInfo/test_2D.cpp b/test_conformance/images/clGetInfo/test_2D.cpp index 764b186d9..12c120267 100644 --- a/test_conformance/images/clGetInfo/test_2D.cpp +++ b/test_conformance/images/clGetInfo/test_2D.cpp @@ -19,6 +19,7 @@ int test_get_image_info_single( cl_context context, image_descriptor *imageInfo, { int error; clMemWrapper image; + clMemWrapper buffer; cl_image_desc imageDesc; void *host_ptr = NULL; @@ -69,6 +70,24 @@ int test_get_image_info_single( cl_context context, image_descriptor *imageInfo, if ( gDebugTrace ) log_info( " - Creating 2D image array %d by %d by %d with flags=0x%lx row_pitch=%d slice_pitch=%d host_ptr=%p...\n", (int)imageInfo->width, (int)imageInfo->height, (int)imageInfo->arraySize, (unsigned long)flags, (int)row_pitch, (int)slice_pitch, host_ptr ); break; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: + if (gDebugTrace) + log_info(" - Creating 1D buffer image %d with flags=0x%lx " + "row_pitch=%d slice_pitch=%d host_ptr=%p...\n", + (int)imageInfo->width, (unsigned long)flags, + (int)row_pitch, (int)slice_pitch, host_ptr); + int err; + buffer = clCreateBuffer(context, flags, imageInfo->rowPitch, + host_ptr, &err); + if (err != CL_SUCCESS) + { + log_error("ERROR: Unable to create buffer for 1D image buffer " + "of size %d (%s)", + (int)imageInfo->rowPitch, IGetErrorString(err)); + return -1; + } + imageDesc.buffer = imageInfo->buffer = buffer; + break; } image = clCreateImage(context, flags, imageInfo->format, &imageDesc, host_ptr, &error); @@ -92,6 +111,11 @@ int test_get_image_info_single( cl_context context, image_descriptor *imageInfo, case CL_MEM_OBJECT_IMAGE2D_ARRAY: log_error( "ERROR: Unable to create 2D image array of size %d x %d x %d (%s)", (int)imageInfo->width, (int)imageInfo->height, (int)imageInfo->arraySize, IGetErrorString( error ) ); break; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: + log_error( + "ERROR: Unable to create 1D image buffer of size %d (%s)", + (int)imageInfo->width, IGetErrorString(error)); + break; } return -1; } @@ -148,6 +172,7 @@ int test_get_image_info_single( cl_context context, image_descriptor *imageInfo, switch (imageInfo->type) { case CL_MEM_OBJECT_IMAGE1D: + case CL_MEM_OBJECT_IMAGE1D_BUFFER: case CL_MEM_OBJECT_IMAGE1D_ARRAY: required_height = 0; break; @@ -175,8 +200,7 @@ int test_get_image_info_single( cl_context context, image_descriptor *imageInfo, case CL_MEM_OBJECT_IMAGE2D: case CL_MEM_OBJECT_IMAGE1D_ARRAY: case CL_MEM_OBJECT_IMAGE2D_ARRAY: - required_depth = 0; - break; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: required_depth = 0; break; case CL_MEM_OBJECT_IMAGE3D: required_depth = imageInfo->depth; break; @@ -198,8 +222,7 @@ int test_get_image_info_single( cl_context context, image_descriptor *imageInfo, case CL_MEM_OBJECT_IMAGE1D: case CL_MEM_OBJECT_IMAGE2D: case CL_MEM_OBJECT_IMAGE3D: - required_array_size = 0; - break; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: required_array_size = 0; break; case CL_MEM_OBJECT_IMAGE1D_ARRAY: case CL_MEM_OBJECT_IMAGE2D_ARRAY: required_array_size = imageInfo->arraySize; diff --git a/test_conformance/images/clGetInfo/test_loops.cpp b/test_conformance/images/clGetInfo/test_loops.cpp index 17f02d8b4..d5c626061 100644 --- a/test_conformance/images/clGetInfo/test_loops.cpp +++ b/test_conformance/images/clGetInfo/test_loops.cpp @@ -21,6 +21,10 @@ extern int test_get_image_info_2D( cl_device_id device, cl_context context, cl_i extern int test_get_image_info_3D( cl_device_id device, cl_context context, cl_image_format *format, cl_mem_flags flags ); extern int test_get_image_info_1D_array( cl_device_id device, cl_context context, cl_image_format *format, cl_mem_flags flags ); extern int test_get_image_info_2D_array( cl_device_id device, cl_context context, cl_image_format *format, cl_mem_flags flags ); +extern int test_get_image_info_1D_buffer(cl_device_id device, + cl_context context, + cl_image_format *format, + cl_mem_flags flags); int test_image_type( cl_device_id device, cl_context context, cl_mem_object_type image_type, cl_mem_flags flags ) { @@ -66,6 +70,10 @@ int test_image_type( cl_device_id device, cl_context context, cl_mem_object_type case CL_MEM_OBJECT_IMAGE2D_ARRAY: test_return = test_get_image_info_2D_array( device, context, &formatList[ i ], flags ); break; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: + test_return = test_get_image_info_1D_buffer( + device, context, &formatList[i], flags); + break; } if (test_return) { diff --git a/test_conformance/images/clReadWriteImage/CMakeLists.txt b/test_conformance/images/clReadWriteImage/CMakeLists.txt index bc1600ff8..9e0bbbc9c 100644 --- a/test_conformance/images/clReadWriteImage/CMakeLists.txt +++ b/test_conformance/images/clReadWriteImage/CMakeLists.txt @@ -3,6 +3,7 @@ set(MODULE_NAME CL_READ_WRITE_IMAGES) set(${MODULE_NAME}_SOURCES main.cpp test_read_1D.cpp + test_read_1D_buffer.cpp test_read_1D_array.cpp test_read_2D.cpp test_read_2D_array.cpp diff --git a/test_conformance/images/clReadWriteImage/main.cpp b/test_conformance/images/clReadWriteImage/main.cpp index 18c7e2397..642b8a572 100644 --- a/test_conformance/images/clReadWriteImage/main.cpp +++ b/test_conformance/images/clReadWriteImage/main.cpp @@ -51,13 +51,15 @@ int test_2Darray(cl_device_id device, cl_context context, cl_command_queue queue { return test_image_set( device, context, queue, CL_MEM_OBJECT_IMAGE2D_ARRAY ); } +int test_1Dbuffer(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return test_image_set(device, context, queue, CL_MEM_OBJECT_IMAGE1D_BUFFER); +} test_definition test_list[] = { - ADD_TEST( 1D ), - ADD_TEST( 2D ), - ADD_TEST( 3D ), - ADD_TEST( 1Darray ), - ADD_TEST( 2Darray ), + ADD_TEST(1D), ADD_TEST(2D), ADD_TEST(3D), + ADD_TEST(1Darray), ADD_TEST(2Darray), ADD_TEST(1Dbuffer), }; const int test_num = ARRAY_SIZE( test_list ); diff --git a/test_conformance/images/clReadWriteImage/test_loops.cpp b/test_conformance/images/clReadWriteImage/test_loops.cpp index 782e4b37d..175a49b8f 100644 --- a/test_conformance/images/clReadWriteImage/test_loops.cpp +++ b/test_conformance/images/clReadWriteImage/test_loops.cpp @@ -33,6 +33,11 @@ extern int test_read_image_set_2D_array(cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, cl_mem_flags flags); +extern int test_read_image_set_1D_buffer(cl_device_id device, + cl_context context, + cl_command_queue queue, + cl_image_format *format, + cl_mem_flags flags); int test_image_type( cl_device_id device, cl_context context, cl_command_queue queue, cl_mem_object_type imageType, cl_mem_flags flags ) { @@ -97,6 +102,10 @@ int test_image_type( cl_device_id device, cl_context context, cl_command_queue q test_return = test_read_image_set_2D_array( device, context, queue, &formatList[i], flags); break; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: + test_return = test_read_image_set_1D_buffer( + device, context, queue, &formatList[i], flags); + break; } if (test_return) diff --git a/test_conformance/images/clReadWriteImage/test_read_1D_buffer.cpp b/test_conformance/images/clReadWriteImage/test_read_1D_buffer.cpp new file mode 100644 index 000000000..01dd80dc3 --- /dev/null +++ b/test_conformance/images/clReadWriteImage/test_read_1D_buffer.cpp @@ -0,0 +1,256 @@ +// +// 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. +// +#include "../testBase.h" +#include + +int test_read_image_1D_buffer(cl_context context, cl_command_queue queue, + image_descriptor *imageInfo, MTdata d, + cl_mem_flags flags) +{ + int error; + + clMemWrapper image; + clMemWrapper buffer; + + // Generate some data to test against + BufferOwningPtr imageValues; + generate_random_image_data(imageInfo, imageValues, d); + + if (gDebugTrace) + { + log_info(" - Creating 1D image %d...\n", (int)imageInfo->width); + log_info(" with %llu mip levels\n", + (unsigned long long)imageInfo->num_mip_levels); + } + + buffer = clCreateBuffer(context, flags, imageInfo->rowPitch, NULL, &error); + if (error != CL_SUCCESS) + { + log_error("ERROR: Unable to create buffer for 1D image buffer of size " + "%d (%s)", + (int)imageInfo->rowPitch, IGetErrorString(error)); + } + + image = create_image_1d(context, flags, imageInfo->format, imageInfo->width, + 0, NULL, buffer, &error); + if (image == NULL) + { + log_error("ERROR: Unable to create 1D image buffer of size %d (%s)", + (int)imageInfo->width, IGetErrorString(error)); + return -1; + } + + if (gDebugTrace) log_info(" - Writing image...\n"); + + size_t origin[3] = { 0, 0, 0 }; + size_t region[3] = { imageInfo->width, 1, 1 }; + size_t fullImageSize = imageInfo->rowPitch; + + BufferOwningPtr resultValues(malloc(fullImageSize)); + size_t imgValMipLevelOffset = 0; + + error = clEnqueueWriteImage(queue, image, CL_FALSE, origin, region, + (gEnablePitch ? imageInfo->rowPitch : 0), 0, + (char *)imageValues + imgValMipLevelOffset, 0, + NULL, NULL); + if (error != CL_SUCCESS) + { + log_error("ERROR: Unable to write to 1D image of size %d \n", + (int)imageInfo->width); + return -1; + } + + // To verify, we just read the results right back and see whether they + // match the input + if (gDebugTrace) + { + log_info(" - Initing result array...\n"); + } + + // Note: we read back without any pitch, to verify pitch actually WORKED + size_t scanlineSize = imageInfo->width * get_pixel_size(imageInfo->format); + size_t imageSize = scanlineSize; + memset(resultValues, 0xff, imageSize); + + if (gDebugTrace) log_info(" - Reading results...\n"); + + error = clEnqueueReadImage(queue, image, CL_TRUE, origin, region, 0, 0, + resultValues, 0, NULL, NULL); + test_error(error, "Unable to read image values"); + + // Verify scanline by scanline, since the pitches are different + char *sourcePtr = (char *)imageValues + imgValMipLevelOffset; + char *destPtr = resultValues; + + if (memcmp(sourcePtr, destPtr, scanlineSize) != 0) + { + log_error("ERROR: Scanline did not verify for image size %d pitch " + "%d (extra %d bytes)\n", + (int)imageInfo->width, (int)imageInfo->rowPitch, + (int)imageInfo->rowPitch + - (int)imageInfo->width + * (int)get_pixel_size(imageInfo->format)); + + log_error("First few values: \n"); + log_error(" Input: "); + uint32_t *s = (uint32_t *)sourcePtr; + uint32_t *d = (uint32_t *)destPtr; + for (int q = 0; q < 12; q++) log_error("%08x ", s[q]); + log_error("\nOutput: "); + for (int q = 0; q < 12; q++) log_error("%08x ", d[q]); + log_error("\n"); + + int outX; + int offset = (int)get_pixel_size(imageInfo->format) + * (int)(imageInfo->width - 16); + if (offset < 0) offset = 0; + int foundCount = debug_find_vector_in_image( + (char *)imageValues + imgValMipLevelOffset, imageInfo, + destPtr + offset, get_pixel_size(imageInfo->format), &outX, NULL, + NULL); + if (foundCount > 0) + { + int returnedOffset = + (offset / (int)get_pixel_size(imageInfo->format)) - outX; + + if (memcmp(sourcePtr + + returnedOffset * get_pixel_size(imageInfo->format), + destPtr, get_pixel_size(imageInfo->format) * 8) + == 0) + log_error(" Values appear to be offsetted by %d\n", + returnedOffset); + else + log_error( + " Calculated offset is %d but unable to verify\n", + returnedOffset); + } + else + { + log_error(" Unable to determine offset\n"); + } + return -1; + } + imgValMipLevelOffset += + imageInfo->width * get_pixel_size(imageInfo->format); + return 0; +} + +int test_read_image_set_1D_buffer(cl_device_id device, cl_context context, + cl_command_queue queue, + cl_image_format *format, cl_mem_flags flags) +{ + size_t maxWidth; + cl_ulong maxAllocSize, memSize; + image_descriptor imageInfo = { 0 }; + RandomSeed seed(gRandomSeed); + size_t pixelSize; + + if (gTestMipmaps) + { + // 1D image buffers don't support mipmaps + // https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_Ext.html#cl_khr_mipmap_image + return 0; + } + + imageInfo.type = CL_MEM_OBJECT_IMAGE1D_BUFFER; + imageInfo.format = format; + imageInfo.height = imageInfo.depth = imageInfo.slicePitch = 0; + pixelSize = get_pixel_size(imageInfo.format); + + int error = clGetDeviceInfo(device, CL_DEVICE_IMAGE_MAX_BUFFER_SIZE, + sizeof(maxWidth), &maxWidth, NULL); + error |= clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, + sizeof(maxAllocSize), &maxAllocSize, NULL); + error |= clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(memSize), + &memSize, NULL); + test_error(error, "Unable to get max image 2D size from device"); + + if (memSize > (cl_ulong)SIZE_MAX) + { + memSize = (cl_ulong)SIZE_MAX; + maxAllocSize = (cl_ulong)SIZE_MAX; + } + + if (gTestSmallImages) + { + for (imageInfo.width = 1; imageInfo.width < 13; imageInfo.width++) + { + imageInfo.rowPitch = imageInfo.width * pixelSize; + + if (gDebugTrace) log_info(" at size %d\n", (int)imageInfo.width); + + int ret = test_read_image_1D_buffer(context, queue, &imageInfo, + seed, flags); + if (ret) return -1; + } + } + else if (gTestMaxImages) + { + // Try a specific set of maximum sizes + size_t numbeOfSizes; + size_t sizes[100][3]; + + get_max_sizes(&numbeOfSizes, 100, sizes, maxWidth, 1, 1, 1, + maxAllocSize, memSize, CL_MEM_OBJECT_IMAGE1D_BUFFER, + imageInfo.format); + + for (size_t idx = 0; idx < numbeOfSizes; idx++) + { + imageInfo.width = sizes[idx][0]; + imageInfo.rowPitch = imageInfo.width * pixelSize; + + log_info("Testing %d\n", (int)imageInfo.width); + if (gDebugTrace) log_info(" at max size %d\n", (int)maxWidth); + if (test_read_image_1D_buffer(context, queue, &imageInfo, seed, + flags)) + return -1; + } + } + else + { + for (int i = 0; i < NUM_IMAGE_ITERATIONS; i++) + { + cl_ulong size; + // Loop until we get a size that a) will fit in the max alloc size + // and b) that an allocation of that image, the result array, plus + // offset arrays, will fit in the global ram space + do + { + imageInfo.width = + (size_t)random_log_in_range(16, (int)maxWidth / 32, seed); + + imageInfo.rowPitch = imageInfo.width * pixelSize; + if (gEnablePitch) + { + size_t extraWidth = (int)random_log_in_range(0, 64, seed); + imageInfo.rowPitch += extraWidth * pixelSize; + } + + size = (size_t)imageInfo.rowPitch * 4; + } while (size > maxAllocSize || (size / 3) > memSize); + + if (gDebugTrace) + log_info(" at size %d (row pitch %d) out of %d\n", + (int)imageInfo.width, (int)imageInfo.rowPitch, + (int)maxWidth); + int ret = test_read_image_1D_buffer(context, queue, &imageInfo, + seed, flags); + if (ret) return -1; + } + } + + return 0; +} diff --git a/test_conformance/images/kernel_image_methods/CMakeLists.txt b/test_conformance/images/kernel_image_methods/CMakeLists.txt index b06e7d5c8..607cbdee0 100644 --- a/test_conformance/images/kernel_image_methods/CMakeLists.txt +++ b/test_conformance/images/kernel_image_methods/CMakeLists.txt @@ -3,6 +3,7 @@ set(MODULE_NAME KERNEL_IMAGE_METHODS) set(${MODULE_NAME}_SOURCES main.cpp test_1D.cpp + test_1D_buffer.cpp test_1D_array.cpp test_2D.cpp test_2D_array.cpp diff --git a/test_conformance/images/kernel_image_methods/main.cpp b/test_conformance/images/kernel_image_methods/main.cpp index 50653ef55..a0a6fc57c 100644 --- a/test_conformance/images/kernel_image_methods/main.cpp +++ b/test_conformance/images/kernel_image_methods/main.cpp @@ -51,13 +51,15 @@ int test_2Darray(cl_device_id device, cl_context context, cl_command_queue queue { return test_image_set( device, context, queue, CL_MEM_OBJECT_IMAGE2D_ARRAY ); } +int test_1Dbuffer(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return test_image_set(device, context, queue, CL_MEM_OBJECT_IMAGE1D_BUFFER); +} test_definition test_list[] = { - ADD_TEST( 1D ), - ADD_TEST( 2D ), - ADD_TEST( 3D ), - ADD_TEST( 1Darray ), - ADD_TEST( 2Darray ), + ADD_TEST(1D), ADD_TEST(2D), ADD_TEST(3D), + ADD_TEST(1Darray), ADD_TEST(2Darray), ADD_TEST(1Dbuffer), }; const int test_num = ARRAY_SIZE( test_list ); diff --git a/test_conformance/images/kernel_image_methods/test_1D_buffer.cpp b/test_conformance/images/kernel_image_methods/test_1D_buffer.cpp new file mode 100644 index 000000000..1885cbd63 --- /dev/null +++ b/test_conformance/images/kernel_image_methods/test_1D_buffer.cpp @@ -0,0 +1,282 @@ +// +// 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. +// +#include "../testBase.h" +#include + + +struct image_kernel_data +{ + cl_int width; + cl_int channelType; + cl_int channelOrder; + cl_int expectedChannelType; + cl_int expectedChannelOrder; +}; + +static const char *methodTest1DImageKernelPattern = + "typedef struct {\n" + " int width;\n" + " int channelType;\n" + " int channelOrder;\n" + " int expectedChannelType;\n" + " int expectedChannelOrder;\n" + " } image_kernel_data;\n" + "__kernel void sample_kernel( %s image1d_buffer_t input, __global " + "image_kernel_data *outData )\n" + "{\n" + " outData->width = get_image_width( input );\n" + " outData->channelType = get_image_channel_data_type( input );\n" + " outData->channelOrder = get_image_channel_order( input );\n" + "\n" + " outData->expectedChannelType = %s;\n" + " outData->expectedChannelOrder = %s;\n" + "}"; + +static int test_get_1Dimage_buffer_info_single(cl_context context, + cl_command_queue queue, + image_descriptor *imageInfo, + MTdata d, cl_mem_flags flags) +{ + int error = 0; + + clProgramWrapper program; + clKernelWrapper kernel; + clMemWrapper image, outDataBuffer, buffer; + char programSrc[10240]; + + image_kernel_data outKernelData; + + // Generate some data to test against + BufferOwningPtr imageValues; + generate_random_image_data(imageInfo, imageValues, d); + + // Construct testing source + if (gDebugTrace) + log_info(" - Creating 1D image %d ...\n", (int)imageInfo->width); + + buffer = clCreateBuffer( + context, flags, imageInfo->width * get_pixel_size(imageInfo->format), + NULL, &error); + if (error != CL_SUCCESS) + { + log_error("ERROR: Unable to create buffer for 1D image buffer of size " + "%d (%s)", + (int)imageInfo->rowPitch, IGetErrorString(error)); + } + + image = create_image_1d(context, flags, imageInfo->format, imageInfo->width, + imageInfo->rowPitch, NULL, buffer, &error); + if (image == NULL) + { + log_error("ERROR: Unable to create 1D image of size %d (%s)", + (int)imageInfo->width, IGetErrorString(error)); + return -1; + } + + char channelTypeConstantString[256] = { 0 }; + char channelOrderConstantString[256] = { 0 }; + + const char *channelTypeName = + GetChannelTypeName(imageInfo->format->image_channel_data_type); + const char *channelOrderName = + GetChannelOrderName(imageInfo->format->image_channel_order); + const char *image_access_qualifier = + (flags == CL_MEM_READ_ONLY) ? "read_only" : "write_only"; + + if (channelTypeName && strlen(channelTypeName)) + sprintf(channelTypeConstantString, "CLK_%s", + &channelTypeName[3]); // replace CL_* with CLK_* + + if (channelOrderName && strlen(channelOrderName)) + sprintf(channelOrderConstantString, "CLK_%s", + &channelOrderName[3]); // replace CL_* with CLK_* + + // Create a program to run against + sprintf(programSrc, methodTest1DImageKernelPattern, image_access_qualifier, + channelTypeConstantString, channelOrderConstantString); + + // log_info("-----------------------------------\n%s\n", programSrc); + error = clFinish(queue); + if (error) print_error(error, "clFinish failed.\n"); + const char *ptr = programSrc; + error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr, + "sample_kernel"); + test_error(error, "Unable to create kernel to test against"); + + // Create an output buffer + outDataBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE, + sizeof(outKernelData), NULL, &error); + test_error(error, "Unable to create output buffer"); + + // Set up arguments and run + error = clSetKernelArg(kernel, 0, sizeof(image), &image); + test_error(error, "Unable to set kernel argument"); + error = clSetKernelArg(kernel, 1, sizeof(outDataBuffer), &outDataBuffer); + test_error(error, "Unable to set kernel argument"); + + size_t threads[1] = { 1 }, localThreads[1] = { 1 }; + + error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, + localThreads, 0, NULL, NULL); + test_error(error, "Unable to run kernel"); + + error = clEnqueueReadBuffer(queue, outDataBuffer, CL_TRUE, 0, + sizeof(outKernelData), &outKernelData, 0, NULL, + NULL); + test_error(error, "Unable to read data buffer"); + + + // Verify the results now + if (outKernelData.width != (cl_int)imageInfo->width) + { + log_error( + "ERROR: Returned width did not validate (expected %d, got %d)\n", + (int)imageInfo->width, (int)outKernelData.width); + error = -1; + } + if (outKernelData.channelType != (cl_int)outKernelData.expectedChannelType) + { + log_error( + "ERROR: Returned channel type did not validate (expected %s (%d), " + "got %d)\n", + GetChannelTypeName(imageInfo->format->image_channel_data_type), + (int)outKernelData.expectedChannelType, + (int)outKernelData.channelType); + error = -1; + } + if (outKernelData.channelOrder + != (cl_int)outKernelData.expectedChannelOrder) + { + log_error("ERROR: Returned channel order did not validate (expected %s " + "(%d), got %d)\n", + GetChannelOrderName(imageInfo->format->image_channel_order), + (int)outKernelData.expectedChannelOrder, + (int)outKernelData.channelOrder); + error = -1; + } + + if (clFinish(queue) != CL_SUCCESS) + { + log_error("ERROR: CL Finished failed in %s \n", __FUNCTION__); + error = -1; + } + + return error; +} + +int test_get_image_info_1D_buffer(cl_device_id device, cl_context context, + cl_command_queue queue, + cl_image_format *format, cl_mem_flags flags) +{ + size_t maxWidth; + cl_ulong maxAllocSize, memSize; + image_descriptor imageInfo = { 0 }; + RandomSeed seed(gRandomSeed); + size_t pixelSize; + + imageInfo.type = CL_MEM_OBJECT_IMAGE1D_BUFFER; + imageInfo.format = format; + imageInfo.height = imageInfo.depth = imageInfo.slicePitch = 0; + pixelSize = get_pixel_size(imageInfo.format); + + int error = clGetDeviceInfo(device, CL_DEVICE_IMAGE_MAX_BUFFER_SIZE, + sizeof(maxWidth), &maxWidth, NULL); + error |= clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, + sizeof(maxAllocSize), &maxAllocSize, NULL); + error |= clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(memSize), + &memSize, NULL); + test_error(error, "Unable to get max image 1D size from device"); + + if (memSize > (cl_ulong)SIZE_MAX) + { + memSize = (cl_ulong)SIZE_MAX; + maxAllocSize = (cl_ulong)SIZE_MAX; + } + + if (gTestSmallImages) + { + for (imageInfo.width = 1; imageInfo.width < 13; imageInfo.width++) + { + imageInfo.rowPitch = imageInfo.width * pixelSize; + if (gDebugTrace) log_info(" at size %d\n", (int)imageInfo.width); + + int ret = test_get_1Dimage_buffer_info_single( + context, queue, &imageInfo, seed, flags); + if (ret) return -1; + } + } + else if (gTestMaxImages) + { + // Try a specific set of maximum sizes + size_t numbeOfSizes; + size_t sizes[100][3]; + + get_max_sizes(&numbeOfSizes, 100, sizes, maxWidth, 1, 1, 1, + maxAllocSize, memSize, CL_MEM_OBJECT_IMAGE1D_BUFFER, + imageInfo.format); + + for (size_t idx = 0; idx < numbeOfSizes; idx++) + { + imageInfo.width = sizes[idx][0]; + imageInfo.rowPitch = imageInfo.width * pixelSize; + + log_info("Testing %d\n", (int)sizes[idx][0]); + if (gDebugTrace) + log_info(" at max size %d\n", (int)sizes[idx][0]); + if (test_get_1Dimage_buffer_info_single(context, queue, &imageInfo, + seed, flags)) + return -1; + } + } + else + { + for (int i = 0; i < NUM_IMAGE_ITERATIONS; i++) + { + cl_ulong size; + // Loop until we get a size that a) will fit in the max alloc size + // and b) that an allocation of that image, the result array, plus + // offset arrays, will fit in the global ram space + do + { + imageInfo.width = + (size_t)random_log_in_range(16, (int)maxWidth / 32, seed); + + imageInfo.rowPitch = imageInfo.width * pixelSize; + size_t extraWidth = (int)random_log_in_range(0, 64, seed); + imageInfo.rowPitch += extraWidth; + + do + { + extraWidth++; + imageInfo.rowPitch += extraWidth; + } while ((imageInfo.rowPitch % pixelSize) != 0); + + size = (cl_ulong)imageInfo.rowPitch * (cl_ulong)imageInfo.height + * 4; + } while (size > maxAllocSize || (size * 3) > memSize); + + if (gDebugTrace) + log_info(" at size %d (row pitch %d) out of %d\n", + (int)imageInfo.width, (int)imageInfo.rowPitch, + (int)maxWidth); + int ret = test_get_1Dimage_buffer_info_single( + context, queue, &imageInfo, seed, flags); + if (ret) return -1; + } + } + + return 0; +} diff --git a/test_conformance/images/kernel_image_methods/test_loops.cpp b/test_conformance/images/kernel_image_methods/test_loops.cpp index 1d892a9b0..c04b1338a 100644 --- a/test_conformance/images/kernel_image_methods/test_loops.cpp +++ b/test_conformance/images/kernel_image_methods/test_loops.cpp @@ -34,6 +34,11 @@ extern int test_get_image_info_2D_array(cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, cl_mem_flags flags); +extern int test_get_image_info_1D_buffer(cl_device_id device, + cl_context context, + cl_command_queue queue, + cl_image_format *format, + cl_mem_flags flags); int test_image_type( cl_device_id device, cl_context context, cl_command_queue queue, cl_mem_object_type imageType, cl_mem_flags flags ) { @@ -84,6 +89,10 @@ int test_image_type( cl_device_id device, cl_context context, cl_command_queue q test_return = test_get_image_info_2D_array( device, context, queue, &formatList[i], flags); break; + case CL_MEM_OBJECT_IMAGE1D_BUFFER: + test_return = test_get_image_info_1D_buffer( + device, context, queue, &formatList[i], flags); + break; } if (test_return) { @@ -109,6 +118,8 @@ int test_image_set( cl_device_id device, cl_context context, cl_command_queue qu switch (imageType) { case CL_MEM_OBJECT_IMAGE1D: test_missing_feature(version_check, "image_1D"); + case CL_MEM_OBJECT_IMAGE1D_BUFFER: + test_missing_feature(version_check, "image_1D_buffer"); case CL_MEM_OBJECT_IMAGE1D_ARRAY: test_missing_feature(version_check, "image_1D_array"); case CL_MEM_OBJECT_IMAGE2D_ARRAY: diff --git a/test_conformance/images/testBase.h b/test_conformance/images/testBase.h index ad48f10d5..2e5eea967 100644 --- a/test_conformance/images/testBase.h +++ b/test_conformance/images/testBase.h @@ -42,17 +42,20 @@ enum TypesToTest // For the clCopyImage test enum MethodsToTest { - k1D = ( 1 << 0 ), - k2D = ( 1 << 1 ), - k1DArray = ( 1 << 2 ), - k2DArray = ( 1 << 3 ), - k3D = ( 1 << 4 ), - k2DTo3D = ( 1 << 5 ), - k3DTo2D = ( 1 << 6 ), - k2DArrayTo2D = ( 1 << 7 ), - k2DTo2DArray = ( 1 << 8 ), - k2DArrayTo3D = ( 1 << 9 ), - k3DTo2DArray = ( 1 << 10 ), + k1D = (1 << 0), + k2D = (1 << 1), + k1DArray = (1 << 2), + k2DArray = (1 << 3), + k3D = (1 << 4), + k2DTo3D = (1 << 5), + k3DTo2D = (1 << 6), + k2DArrayTo2D = (1 << 7), + k2DTo2DArray = (1 << 8), + k2DArrayTo3D = (1 << 9), + k3DTo2DArray = (1 << 10), + k1DBuffer = (1 << 11), + k1DTo1DBuffer = (1 << 12), + k1DBufferTo1D = (1 << 13), }; From 61a220ff332e838088be1fc9cf1d124d37c05920 Mon Sep 17 00:00:00 2001 From: David Tobolik <124600480+dtobolik@users.noreply.github.com> Date: Tue, 16 Apr 2024 17:49:24 +0200 Subject: [PATCH 4/7] fix: move suggested_local_size to test_api (#1916) --- test_conformance/api/CMakeLists.txt | 1 + test_conformance/api/main.cpp | 4 + test_conformance/api/procs.h | 15 +++- .../test_wg_suggested_local_work_size.cpp | 2 +- test_conformance/workgroups/CMakeLists.txt | 1 - test_conformance/workgroups/main.cpp | 3 - test_conformance/workgroups/procs.h | 78 ++++++++++++------- 7 files changed, 70 insertions(+), 34 deletions(-) rename test_conformance/{workgroups => api}/test_wg_suggested_local_work_size.cpp (99%) diff --git a/test_conformance/api/CMakeLists.txt b/test_conformance/api/CMakeLists.txt index d3e6c6a7c..5b1f491ce 100644 --- a/test_conformance/api/CMakeLists.txt +++ b/test_conformance/api/CMakeLists.txt @@ -35,6 +35,7 @@ set(${MODULE_NAME}_SOURCES test_mem_object_properties_queries.cpp test_queue_properties_queries.cpp test_pipe_properties_queries.cpp + test_wg_suggested_local_work_size.cpp ) include(../CMakeCommon.txt) diff --git a/test_conformance/api/main.cpp b/test_conformance/api/main.cpp index fa76a4064..e0900df07 100644 --- a/test_conformance/api/main.cpp +++ b/test_conformance/api/main.cpp @@ -148,6 +148,10 @@ test_definition test_list[] = { ADD_TEST(min_image_formats), ADD_TEST(negative_get_platform_info), ADD_TEST(negative_get_platform_ids), + + ADD_TEST(work_group_suggested_local_size_1D), + ADD_TEST(work_group_suggested_local_size_2D), + ADD_TEST(work_group_suggested_local_size_3D), }; const int test_num = ARRAY_SIZE(test_list); diff --git a/test_conformance/api/procs.h b/test_conformance/api/procs.h index 1bcb31162..320ad65a1 100644 --- a/test_conformance/api/procs.h +++ b/test_conformance/api/procs.h @@ -1,6 +1,6 @@ // // Copyright (c) 2017 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 @@ -204,3 +204,16 @@ extern int test_negative_get_platform_ids(cl_device_id deviceID, int num_elements); extern int test_kernel_attributes(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); + +extern int test_work_group_suggested_local_size_1D(cl_device_id device, + cl_context context, + cl_command_queue queue, + int n_elems); +extern int test_work_group_suggested_local_size_2D(cl_device_id device, + cl_context context, + cl_command_queue queue, + int n_elems); +extern int test_work_group_suggested_local_size_3D(cl_device_id device, + cl_context context, + cl_command_queue queue, + int n_elems); diff --git a/test_conformance/workgroups/test_wg_suggested_local_work_size.cpp b/test_conformance/api/test_wg_suggested_local_work_size.cpp similarity index 99% rename from test_conformance/workgroups/test_wg_suggested_local_work_size.cpp rename to test_conformance/api/test_wg_suggested_local_work_size.cpp index a31fca63f..2b2a5404f 100644 --- a/test_conformance/workgroups/test_wg_suggested_local_work_size.cpp +++ b/test_conformance/api/test_wg_suggested_local_work_size.cpp @@ -72,7 +72,7 @@ const char* wg_scan_local_work_group_size = R"( global uint *output) { __local char c[LOCAL_MEM_SIZE]; - + if(!is_zero_linear_id()) return; for (uint i = 0; i < 3; i++) { diff --git a/test_conformance/workgroups/CMakeLists.txt b/test_conformance/workgroups/CMakeLists.txt index 0c004b320..f9f9bad1e 100644 --- a/test_conformance/workgroups/CMakeLists.txt +++ b/test_conformance/workgroups/CMakeLists.txt @@ -6,7 +6,6 @@ set(${MODULE_NAME}_SOURCES test_wg_any.cpp test_wg_broadcast.cpp test_wg_scan_reduce.cpp - test_wg_suggested_local_work_size.cpp ) include(../CMakeCommon.txt) diff --git a/test_conformance/workgroups/main.cpp b/test_conformance/workgroups/main.cpp index abb1145b3..11849e1f6 100644 --- a/test_conformance/workgroups/main.cpp +++ b/test_conformance/workgroups/main.cpp @@ -38,9 +38,6 @@ test_definition test_list[] = { ADD_TEST_VERSION(work_group_broadcast_1D, Version(2, 0)), ADD_TEST_VERSION(work_group_broadcast_2D, Version(2, 0)), ADD_TEST_VERSION(work_group_broadcast_3D, Version(2, 0)), - ADD_TEST(work_group_suggested_local_size_1D), - ADD_TEST(work_group_suggested_local_size_2D), - ADD_TEST(work_group_suggested_local_size_3D) }; const int test_num = ARRAY_SIZE(test_list); diff --git a/test_conformance/workgroups/procs.h b/test_conformance/workgroups/procs.h index 6143d5253..0baa9066a 100644 --- a/test_conformance/workgroups/procs.h +++ b/test_conformance/workgroups/procs.h @@ -20,33 +20,55 @@ #include "harness/conversions.h" #include "harness/mt19937.h" -extern int create_program_and_kernel(const char *source, const char *kernel_name, cl_program *program_ret, cl_kernel *kernel_ret); +extern int create_program_and_kernel(const char *source, + const char *kernel_name, + cl_program *program_ret, + cl_kernel *kernel_ret); -extern int test_work_group_all(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_work_group_any(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_work_group_broadcast_1D(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_work_group_broadcast_2D(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_work_group_broadcast_3D(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_work_group_reduce_add(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_work_group_reduce_min(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_work_group_reduce_max(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +extern int test_work_group_all(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_work_group_any(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_work_group_broadcast_1D(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_work_group_broadcast_2D(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_work_group_broadcast_3D(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_work_group_reduce_add(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_work_group_reduce_min(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_work_group_reduce_max(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); -extern int test_work_group_scan_exclusive_add(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_work_group_scan_exclusive_min(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_work_group_scan_exclusive_max(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_work_group_scan_inclusive_add(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_work_group_scan_inclusive_min(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_work_group_scan_inclusive_max(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); - -extern int test_work_group_suggested_local_size_1D(cl_device_id device, - cl_context context, - cl_command_queue queue, - int n_elems); -extern int test_work_group_suggested_local_size_2D(cl_device_id device, - cl_context context, - cl_command_queue queue, - int n_elems); -extern int test_work_group_suggested_local_size_3D(cl_device_id device, - cl_context context, - cl_command_queue queue, - int n_elems); +extern int test_work_group_scan_exclusive_add(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_work_group_scan_exclusive_min(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_work_group_scan_exclusive_max(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_work_group_scan_inclusive_add(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_work_group_scan_inclusive_min(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); +extern int test_work_group_scan_inclusive_max(cl_device_id deviceID, + cl_context context, + cl_command_queue queue, + int num_elements); From a977e35e0b779f7ebf29688a3e7b29e47fa7910d Mon Sep 17 00:00:00 2001 From: Pedro Olsen Ferreira Date: Tue, 16 Apr 2024 16:53:59 +0100 Subject: [PATCH 5/7] Delete stale files (#1945) These aren't included or built anywhere. --- test_common/autotest/autotest.hpp | 38 ----------------- test_common/autotest/test_suite.hpp | 63 ----------------------------- 2 files changed, 101 deletions(-) delete mode 100644 test_common/autotest/autotest.hpp delete mode 100644 test_common/autotest/test_suite.hpp diff --git a/test_common/autotest/autotest.hpp b/test_common/autotest/autotest.hpp deleted file mode 100644 index a3d14dc02..000000000 --- a/test_common/autotest/autotest.hpp +++ /dev/null @@ -1,38 +0,0 @@ -// -// Copyright (c) 2017 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 TEST_COMMON_AUTOTEST_AUTOTEST_HPP -#define TEST_COMMON_AUTOTEST_AUTOTEST_HPP - -#include "test_suite.hpp" - -#define STR_JOIN( X, Y ) STR_DO_JOIN( X, Y ) -#define STR_DO_JOIN( X, Y ) STR_DO_JOIN_2(X,Y) -#define STR_DO_JOIN_2( X, Y ) X##Y - - -// How to use AUTO_TEST_CASE macro: -// -// AUTO_TEST_CASE()(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems) -// { -// (test case code...) -// } -// -#define AUTO_TEST_CASE(name) \ - struct name { static int run_test(cl_device_id, cl_context, cl_command_queue, int); }; \ - static autotest::detail::test_case_registration STR_JOIN(name, STR_JOIN(_registration, __LINE__)) (#name, name::run_test); \ - int name::run_test - -#endif //TEST_COMMON_AUTOTEST_AUTOTEST_HPP \ No newline at end of file diff --git a/test_common/autotest/test_suite.hpp b/test_common/autotest/test_suite.hpp deleted file mode 100644 index b831b9e64..000000000 --- a/test_common/autotest/test_suite.hpp +++ /dev/null @@ -1,63 +0,0 @@ -// -// Copyright (c) 2017 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 TEST_COMMON_AUTOTEST_TEST_SUITE_HPP -#define TEST_COMMON_AUTOTEST_TEST_SUITE_HPP - -#include -#include - -namespace autotest { - -struct test_suite { - test_suite(const std::string& name) - : name(name) - { - - } - - void add(const test_definition& td) - { - test_defs.push_back(td); - } - - // List of test definitions - std::vector test_defs; - // Test suite name - const std::string name; - - static test_suite& global_test_suite() - { - static test_suite global_test_suite("global"); - return global_test_suite; - } -}; - -namespace detail { - -struct test_case_registration -{ - test_case_registration(const std::string& name, - const test_function_pointer ptr) - { - ::autotest::test_suite::global_test_suite().add( - test_definition({ ptr, strdup(name.c_str()) })); - } -}; - -} // end detail namespace -} // end autotest namespace - -#endif // TEST_COMMON_AUTOTEST_TEST_SUITE_HPP From 43237f74a8609ba95494a606e9a28b7780ffa4ed Mon Sep 17 00:00:00 2001 From: Pedro Olsen Ferreira Date: Tue, 16 Apr 2024 16:54:28 +0100 Subject: [PATCH 6/7] Remove dead code (#1946) This is 7-year old CMake code that never did anything because the function invocation is missing the PROPERTIES keyword. Adding the keyword results in build errors, so just drop the dead code since everything seems to be working without it anyway. --- test_conformance/conversions/CMakeLists.txt | 8 -------- 1 file changed, 8 deletions(-) diff --git a/test_conformance/conversions/CMakeLists.txt b/test_conformance/conversions/CMakeLists.txt index 11106439b..e2e976672 100644 --- a/test_conformance/conversions/CMakeLists.txt +++ b/test_conformance/conversions/CMakeLists.txt @@ -8,14 +8,6 @@ if("${CLConform_TARGET_ARCH}" STREQUAL "ARM" OR "${CLConform_TARGET_ARCH}" STREQ list(APPEND ${MODULE_NAME}_SOURCES fplib.cpp) endif() -if(NOT CMAKE_CL_64 AND NOT MSVC AND NOT ANDROID) -# -march is needed for CPU atomics, default arch on gcc is i386 -# that does not support atomics. -set_source_files_properties( - ${MODULE_NAME}_SOURCES - COMPILE_FLAGS -march=i686) -endif(NOT CMAKE_CL_64 AND NOT MSVC AND NOT ANDROID) - set_gnulike_module_compile_flags("-Wno-sign-compare") include(../CMakeCommon.txt) From 7dc9593d8828157dbce35c83e75bd8999d12dfea Mon Sep 17 00:00:00 2001 From: Sven van Haastregt Date: Tue, 16 Apr 2024 19:11:08 +0200 Subject: [PATCH 7/7] test_vulkan: don't throw from destructors (#1947) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Only report an error (and include the error code), but don't throw an exception as that would call `terminate`. Failure to release resources is not fatal in other parts of the CTS. This fixes `-Wterminate` warnings: warning: ‘throw’ will always call ‘terminate’ [-Wterminate] note: in C++11 destructors default to ‘noexcept’ Signed-off-by: Sven van Haastregt --- .../common/vulkan_wrapper/opencl_vulkan_wrapper.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/test_conformance/common/vulkan_wrapper/opencl_vulkan_wrapper.cpp b/test_conformance/common/vulkan_wrapper/opencl_vulkan_wrapper.cpp index c13758022..f295387a1 100644 --- a/test_conformance/common/vulkan_wrapper/opencl_vulkan_wrapper.cpp +++ b/test_conformance/common/vulkan_wrapper/opencl_vulkan_wrapper.cpp @@ -853,7 +853,7 @@ clExternalImportableSemaphore::~clExternalImportableSemaphore() cl_int err = clReleaseSemaphoreKHRptr(m_externalSemaphore); if (err != CL_SUCCESS) { - throw std::runtime_error("clReleaseSemaphoreKHR failed!"); + log_error("clReleaseSemaphoreKHR failed with %d\n", err); } } @@ -935,7 +935,7 @@ clExternalExportableSemaphore::~clExternalExportableSemaphore() cl_int err = clReleaseSemaphoreKHRptr(m_externalSemaphore); if (err != CL_SUCCESS) { - throw std::runtime_error("clReleaseSemaphoreKHR failed!"); + log_error("clReleaseSemaphoreKHR failed with %d\n", err); } } @@ -1052,4 +1052,4 @@ VulkanImageTiling vkClExternalMemoryHandleTilingAssumption( } return mode; -} \ No newline at end of file +}