From 7d4e5827e2adb16ff981344e148284796fb97f99 Mon Sep 17 00:00:00 2001 From: AhmedAmraniAkdi Date: Wed, 1 May 2024 16:41:17 +0100 Subject: [PATCH 1/7] Add cl_khr_kernel_clock tests --- test_conformance/extensions/CMakeLists.txt | 1 + .../cl_khr_kernel_clock/CMakeLists.txt | 8 + .../cl_khr_kernel_clock/kernel_clock.cpp | 201 ++++++++++++++++++ .../extensions/cl_khr_kernel_clock/main.cpp | 29 +++ .../extensions/cl_khr_kernel_clock/procs.h | 27 +++ 5 files changed, 266 insertions(+) create mode 100644 test_conformance/extensions/cl_khr_kernel_clock/CMakeLists.txt create mode 100644 test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp create mode 100644 test_conformance/extensions/cl_khr_kernel_clock/main.cpp create mode 100644 test_conformance/extensions/cl_khr_kernel_clock/procs.h diff --git a/test_conformance/extensions/CMakeLists.txt b/test_conformance/extensions/CMakeLists.txt index 3187174f22..77aa2f18eb 100644 --- a/test_conformance/extensions/CMakeLists.txt +++ b/test_conformance/extensions/CMakeLists.txt @@ -2,6 +2,7 @@ add_subdirectory( cl_ext_cxx_for_opencl ) add_subdirectory( cl_khr_command_buffer ) add_subdirectory( cl_khr_dx9_media_sharing ) add_subdirectory( cl_khr_semaphore ) +add_subdirectory( cl_khr_kernel_clock ) if(VULKAN_IS_SUPPORTED) add_subdirectory( cl_khr_external_semaphore ) endif() diff --git a/test_conformance/extensions/cl_khr_kernel_clock/CMakeLists.txt b/test_conformance/extensions/cl_khr_kernel_clock/CMakeLists.txt new file mode 100644 index 0000000000..b0fe19375b --- /dev/null +++ b/test_conformance/extensions/cl_khr_kernel_clock/CMakeLists.txt @@ -0,0 +1,8 @@ +set(MODULE_NAME CL_KHR_KERNEL_CLOCK) + +set(${MODULE_NAME}_SOURCES + main.cpp + kernel_clock.cpp +) + +include(../../CMakeCommon.txt) \ No newline at end of file diff --git a/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp b/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp new file mode 100644 index 0000000000..9867d651a2 --- /dev/null +++ b/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp @@ -0,0 +1,201 @@ +// 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 "procs.h" +#include "harness/typeWrappers.h" + +// write 1 to the output if the clock did not increase +static const char *kernel_sources[2] = { + { + "__kernel void SampleClock(__global uint* buf)\n" + "{\n" + " ulong time1, time2;\n" + " time1 = clock_read_%s();\n" + " time2 = clock_read_%s();\n" + " if(time1 > time2)\n" + " {\n" + " buf[0] = 1;\n" + " }\n" + "}\n", + }, + { + "__kernel void SampleClock(__global uint* buf)\n" + "{\n" + " uint2 time1, time2;\n" + " time1 = clock_read_hilo_%s();\n" + " time2 = clock_read_hilo_%s();\n" + " if (time1.hi > time2.hi || (time1.hi == time2.hi && time1.lo > " + "time2.lo))\n" + " {\n" + " buf[0] = 1;\n" + " }\n" + "}\n", + } +}; + +class KernelClockTest { + +public: + KernelClockTest(cl_device_id device, cl_context context, + cl_command_queue queue, + cl_kernel_clock_capabilities capability) + : device(device), context(context), queue(queue), capability(capability) + {} + + bool Skip() + { + cl_kernel_clock_capabilities capabilities; + cl_int error = clGetDeviceInfo( + device, CL_DEVICE_KERNEL_CLOCK_CAPABILITIES_KHR, + sizeof(cl_kernel_clock_capabilities), &capabilities, NULL); + test_error(error, + "Unable to query " + "CL_DEVICE_KERNEL_CLOCK_CAPABILITIES_KHR"); + + // Skip if capability is not supported + return capability != (capabilities & capability); + } + + cl_int RunTest() + { + size_t global_size = 1; + cl_uint buf = 0; + char kernel_src[512]; + const char *ptr; + cl_int error; + + // 2 built-ins for each scope + for (size_t i = 0; i < 2; i++) + { + buf = 0; + clProgramWrapper program; + clKernelWrapper kernel; + clMemWrapper out_mem; + + if (i == 0 && !gHasLong) + { + log_info("The device does not support ulong. Testing hilo " + "built-ins only\n"); + continue; + } + + switch (capability) + { + case CL_DEVICE_KERNEL_CLOCK_SCOPE_DEVICE_KHR: { + sprintf(kernel_src, kernel_sources[i], "device", "device"); + break; + } + case CL_DEVICE_KERNEL_CLOCK_SCOPE_WORK_GROUP_KHR: { + sprintf(kernel_src, kernel_sources[i], "workgroup", + "workgroup"); + break; + } + case CL_DEVICE_KERNEL_CLOCK_SCOPE_SUB_GROUP_KHR: { + sprintf(kernel_src, kernel_sources[i], "subgroup", + "subgroup"); + break; + } + } + + ptr = kernel_src; + + error = create_single_kernel_helper_create_program( + context, &program, 1, &ptr); + test_error(error, "Failed to create program with source"); + + error = + clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr); + test_error(error, "Failed to build program"); + + out_mem = clCreateBuffer(context, CL_MEM_WRITE_ONLY, + sizeof(cl_uint), nullptr, &error); + test_error(error, "clCreateBuffer failed"); + + kernel = clCreateKernel(program, "SampleClock", &error); + test_error(error, "Failed to create kernel"); + + error = clSetKernelArg(kernel, 0, sizeof(out_mem), &out_mem); + test_error(error, "clSetKernelArg failed"); + + error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size, + NULL, 0, NULL, NULL); + test_error(error, "clNDRangeKernel failed"); + + error = clEnqueueReadBuffer(queue, out_mem, CL_BLOCKING, 0, + sizeof(cl_uint), &buf, 0, NULL, NULL); + test_error(error, "clEnqueueReadBuffer failed"); + + if (buf == 1) + { + log_error( + "Sampling the clock returned bad values, time1 > time2.\n"); + return TEST_FAIL; + } + } + + return CL_SUCCESS; + } + +private: + cl_device_id device; + cl_context context; + cl_command_queue queue; + cl_kernel_clock_capabilities capability; +}; + +int MakeAndRunTest(cl_device_id device, cl_context context, + cl_command_queue queue, + cl_kernel_clock_capabilities capability) +{ + if (!is_extension_available(device, "cl_khr_kernel_clock")) + { + log_info( + "The device does not support the cl_khr_kernel_clock extension.\n"); + return TEST_SKIPPED_ITSELF; + } + + KernelClockTest test_fixture = + KernelClockTest(device, context, queue, capability); + + if (test_fixture.Skip()) + { + return TEST_SKIPPED_ITSELF; + } + + cl_int error = test_fixture.RunTest(); + test_error_ret(error, "Test Failed", TEST_FAIL); + + return TEST_PASS; +} + +int test_device_scope(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return MakeAndRunTest(device, context, queue, + CL_DEVICE_KERNEL_CLOCK_SCOPE_DEVICE_KHR); +} + +int test_workgroup_scope(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return MakeAndRunTest(device, context, queue, + CL_DEVICE_KERNEL_CLOCK_SCOPE_WORK_GROUP_KHR); +} + +int test_subgroup_scope(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + return MakeAndRunTest(device, context, queue, + CL_DEVICE_KERNEL_CLOCK_SCOPE_SUB_GROUP_KHR); +} diff --git a/test_conformance/extensions/cl_khr_kernel_clock/main.cpp b/test_conformance/extensions/cl_khr_kernel_clock/main.cpp new file mode 100644 index 0000000000..8a2d98554b --- /dev/null +++ b/test_conformance/extensions/cl_khr_kernel_clock/main.cpp @@ -0,0 +1,29 @@ +// 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 "procs.h" +#include "harness/testHarness.h" + +test_definition test_list[] = { + ADD_TEST(device_scope), + ADD_TEST(workgroup_scope), + ADD_TEST(subgroup_scope), +}; + + +int main(int argc, const char *argv[]) +{ + return runTestHarness(argc, argv, ARRAY_SIZE(test_list), test_list, false, + 0); +} diff --git a/test_conformance/extensions/cl_khr_kernel_clock/procs.h b/test_conformance/extensions/cl_khr_kernel_clock/procs.h new file mode 100644 index 0000000000..8411b820ac --- /dev/null +++ b/test_conformance/extensions/cl_khr_kernel_clock/procs.h @@ -0,0 +1,27 @@ +// 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. +// +#ifndef _CL_KHR_KERNEL_CLOCK_PROCS_H +#define _CL_KHR_KERNEL_CLOCK_PROCS_H + +#include + +extern int test_device_scope(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_workgroup_scope(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_subgroup_scope(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements); + +#endif /*_CL_KHR_KERNEL_CLOCK_PROCS_H*/ From 17bb73bf1ec1e3489eee083dacdea33a6cf1526f Mon Sep 17 00:00:00 2001 From: AhmedAmraniAkdi Date: Wed, 1 May 2024 17:15:52 +0100 Subject: [PATCH 2/7] Change cl_kernel_clock_capabilities_khr to cl_device_kernel_clock_capabilities_khr --- .../extensions/cl_khr_kernel_clock/kernel_clock.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp b/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp index 9867d651a2..c3759c1b4b 100644 --- a/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp +++ b/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp @@ -49,16 +49,16 @@ class KernelClockTest { public: KernelClockTest(cl_device_id device, cl_context context, cl_command_queue queue, - cl_kernel_clock_capabilities capability) + cl_device_kernel_clock_capabilities_khr capability) : device(device), context(context), queue(queue), capability(capability) {} bool Skip() { - cl_kernel_clock_capabilities capabilities; + cl_device_kernel_clock_capabilities_khr capabilities; cl_int error = clGetDeviceInfo( device, CL_DEVICE_KERNEL_CLOCK_CAPABILITIES_KHR, - sizeof(cl_kernel_clock_capabilities), &capabilities, NULL); + sizeof(cl_device_kernel_clock_capabilities_khr), &capabilities, NULL); test_error(error, "Unable to query " "CL_DEVICE_KERNEL_CLOCK_CAPABILITIES_KHR"); @@ -151,12 +151,12 @@ class KernelClockTest { cl_device_id device; cl_context context; cl_command_queue queue; - cl_kernel_clock_capabilities capability; + cl_device_kernel_clock_capabilities_khr capability; }; int MakeAndRunTest(cl_device_id device, cl_context context, cl_command_queue queue, - cl_kernel_clock_capabilities capability) + cl_device_kernel_clock_capabilities_khr capability) { if (!is_extension_available(device, "cl_khr_kernel_clock")) { From 3b96da891e9d306186595ce5c4268dec7d289880 Mon Sep 17 00:00:00 2001 From: AhmedAmraniAkdi Date: Wed, 1 May 2024 17:17:55 +0100 Subject: [PATCH 3/7] Formatting --- .../extensions/cl_khr_kernel_clock/kernel_clock.cpp | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp b/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp index c3759c1b4b..8d5211a0f8 100644 --- a/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp +++ b/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp @@ -56,9 +56,10 @@ class KernelClockTest { bool Skip() { cl_device_kernel_clock_capabilities_khr capabilities; - cl_int error = clGetDeviceInfo( - device, CL_DEVICE_KERNEL_CLOCK_CAPABILITIES_KHR, - sizeof(cl_device_kernel_clock_capabilities_khr), &capabilities, NULL); + cl_int error = + clGetDeviceInfo(device, CL_DEVICE_KERNEL_CLOCK_CAPABILITIES_KHR, + sizeof(cl_device_kernel_clock_capabilities_khr), + &capabilities, NULL); test_error(error, "Unable to query " "CL_DEVICE_KERNEL_CLOCK_CAPABILITIES_KHR"); From 6ce88df3fd17b88b6c76c99f9f267616ca538111 Mon Sep 17 00:00:00 2001 From: AhmedAmraniAkdi Date: Wed, 1 May 2024 17:20:32 +0100 Subject: [PATCH 4/7] Add cl_khr_kernel_clock to the list of known extensions in the compiler defines for extension subtest --- .../compiler/test_compiler_defines_for_extensions.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/test_conformance/compiler/test_compiler_defines_for_extensions.cpp b/test_conformance/compiler/test_compiler_defines_for_extensions.cpp index 623c2dbfb7..dbbb8c02c2 100644 --- a/test_conformance/compiler/test_compiler_defines_for_extensions.cpp +++ b/test_conformance/compiler/test_compiler_defines_for_extensions.cpp @@ -91,6 +91,7 @@ const char *known_extensions[] = { "cl_khr_external_memory_dma_buf", "cl_khr_command_buffer", "cl_khr_command_buffer_mutable_dispatch", + "cl_khr_kernel_clock", }; size_t num_known_extensions = ARRAY_SIZE(known_extensions); From 3912bfc201ff2ec51cce9b13a5d253a4d746cfcc Mon Sep 17 00:00:00 2001 From: AhmedAmraniAkdi Date: Wed, 1 May 2024 17:25:28 +0100 Subject: [PATCH 5/7] Fix error: braces around scalar initializer --- .../cl_khr_kernel_clock/kernel_clock.cpp | 46 +++++++++---------- 1 file changed, 21 insertions(+), 25 deletions(-) diff --git a/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp b/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp index 8d5211a0f8..abc820256e 100644 --- a/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp +++ b/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp @@ -17,31 +17,27 @@ // write 1 to the output if the clock did not increase static const char *kernel_sources[2] = { - { - "__kernel void SampleClock(__global uint* buf)\n" - "{\n" - " ulong time1, time2;\n" - " time1 = clock_read_%s();\n" - " time2 = clock_read_%s();\n" - " if(time1 > time2)\n" - " {\n" - " buf[0] = 1;\n" - " }\n" - "}\n", - }, - { - "__kernel void SampleClock(__global uint* buf)\n" - "{\n" - " uint2 time1, time2;\n" - " time1 = clock_read_hilo_%s();\n" - " time2 = clock_read_hilo_%s();\n" - " if (time1.hi > time2.hi || (time1.hi == time2.hi && time1.lo > " - "time2.lo))\n" - " {\n" - " buf[0] = 1;\n" - " }\n" - "}\n", - } + "__kernel void SampleClock(__global uint* buf)\n" + "{\n" + " ulong time1, time2;\n" + " time1 = clock_read_%s();\n" + " time2 = clock_read_%s();\n" + " if(time1 > time2)\n" + " {\n" + " buf[0] = 1;\n" + " }\n" + "}\n", + "__kernel void SampleClock(__global uint* buf)\n" + "{\n" + " uint2 time1, time2;\n" + " time1 = clock_read_hilo_%s();\n" + " time2 = clock_read_hilo_%s();\n" + " if (time1.hi > time2.hi || (time1.hi == time2.hi && time1.lo > " + "time2.lo))\n" + " {\n" + " buf[0] = 1;\n" + " }\n" + "}\n", }; class KernelClockTest { From 72c91b4e63113cfe4b42f1f88668019ed4f7ca6f Mon Sep 17 00:00:00 2001 From: AhmedAmraniAkdi Date: Fri, 3 May 2024 17:27:46 +0100 Subject: [PATCH 6/7] Put cl_khr_kernel_clock before the API-only extensions in the known_extensions array. Add newline to CMakeLists.txt Wrap everything under an anonymous namespace, remove extern keyword, use c++ raw strings for the ocl kernels, and remove underscore from CL_KHR_KERNEL_CLOCK_PROCS_H --- .../test_compiler_defines_for_extensions.cpp | 4 +- .../cl_khr_kernel_clock/CMakeLists.txt | 2 +- .../cl_khr_kernel_clock/kernel_clock.cpp | 46 ++++++++++--------- .../extensions/cl_khr_kernel_clock/procs.h | 18 ++++---- 4 files changed, 37 insertions(+), 33 deletions(-) diff --git a/test_conformance/compiler/test_compiler_defines_for_extensions.cpp b/test_conformance/compiler/test_compiler_defines_for_extensions.cpp index dbbb8c02c2..ffd0d6a099 100644 --- a/test_conformance/compiler/test_compiler_defines_for_extensions.cpp +++ b/test_conformance/compiler/test_compiler_defines_for_extensions.cpp @@ -53,6 +53,7 @@ const char *known_extensions[] = { "cl_khr_extended_bit_ops", "cl_khr_integer_dot_product", "cl_khr_subgroup_rotate", + "cl_khr_kernel_clock", // API-only extensions after this point. If you add above here, modify // first_API_extension below. "cl_khr_icd", @@ -91,11 +92,10 @@ const char *known_extensions[] = { "cl_khr_external_memory_dma_buf", "cl_khr_command_buffer", "cl_khr_command_buffer_mutable_dispatch", - "cl_khr_kernel_clock", }; size_t num_known_extensions = ARRAY_SIZE(known_extensions); -size_t first_API_extension = 31; +size_t first_API_extension = 32; const char *known_embedded_extensions[] = { "cles_khr_int64", diff --git a/test_conformance/extensions/cl_khr_kernel_clock/CMakeLists.txt b/test_conformance/extensions/cl_khr_kernel_clock/CMakeLists.txt index b0fe19375b..066ebb65d9 100644 --- a/test_conformance/extensions/cl_khr_kernel_clock/CMakeLists.txt +++ b/test_conformance/extensions/cl_khr_kernel_clock/CMakeLists.txt @@ -5,4 +5,4 @@ set(${MODULE_NAME}_SOURCES kernel_clock.cpp ) -include(../../CMakeCommon.txt) \ No newline at end of file +include(../../CMakeCommon.txt) diff --git a/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp b/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp index abc820256e..d8298b3d63 100644 --- a/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp +++ b/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp @@ -15,29 +15,31 @@ #include "procs.h" #include "harness/typeWrappers.h" +namespace { + // write 1 to the output if the clock did not increase static const char *kernel_sources[2] = { - "__kernel void SampleClock(__global uint* buf)\n" - "{\n" - " ulong time1, time2;\n" - " time1 = clock_read_%s();\n" - " time2 = clock_read_%s();\n" - " if(time1 > time2)\n" - " {\n" - " buf[0] = 1;\n" - " }\n" - "}\n", - "__kernel void SampleClock(__global uint* buf)\n" - "{\n" - " uint2 time1, time2;\n" - " time1 = clock_read_hilo_%s();\n" - " time2 = clock_read_hilo_%s();\n" - " if (time1.hi > time2.hi || (time1.hi == time2.hi && time1.lo > " - "time2.lo))\n" - " {\n" - " buf[0] = 1;\n" - " }\n" - "}\n", + R"(__kernel void SampleClock(__global uint* buf) + { + ulong time1, time2; + time1 = clock_read_%s(); + time2 = clock_read_%s(); + if(time1 > time2) + { + buf[0] = 1; + } + })", + R"(__kernel void SampleClock(__global uint* buf) + { + uint2 time1, time2; + time1 = clock_read_hilo_%s(); + time2 = clock_read_hilo_%s(); + if(time1.hi > time2.hi || (time1.hi == time2.hi && time1.lo > + time2.lo)) + { + buf[0] = 1; + } + })", }; class KernelClockTest { @@ -176,6 +178,8 @@ int MakeAndRunTest(cl_device_id device, cl_context context, return TEST_PASS; } +} + int test_device_scope(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) { diff --git a/test_conformance/extensions/cl_khr_kernel_clock/procs.h b/test_conformance/extensions/cl_khr_kernel_clock/procs.h index 8411b820ac..a82564bc81 100644 --- a/test_conformance/extensions/cl_khr_kernel_clock/procs.h +++ b/test_conformance/extensions/cl_khr_kernel_clock/procs.h @@ -12,16 +12,16 @@ // See the License for the specific language governing permissions and // limitations under the License. // -#ifndef _CL_KHR_KERNEL_CLOCK_PROCS_H -#define _CL_KHR_KERNEL_CLOCK_PROCS_H +#ifndef CL_KHR_KERNEL_CLOCK_PROCS_H +#define CL_KHR_KERNEL_CLOCK_PROCS_H #include -extern int test_device_scope(cl_device_id device, cl_context context, - cl_command_queue queue, int num_elements); -extern int test_workgroup_scope(cl_device_id device, cl_context context, - cl_command_queue queue, int num_elements); -extern int test_subgroup_scope(cl_device_id device, cl_context context, - cl_command_queue queue, int num_elements); +int test_device_scope(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements); +int test_workgroup_scope(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements); +int test_subgroup_scope(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements); -#endif /*_CL_KHR_KERNEL_CLOCK_PROCS_H*/ +#endif /*CL_KHR_KERNEL_CLOCK_PROCS_H*/ From 8b4a2393245c573d7437b7fe9371b365b8a4878a Mon Sep 17 00:00:00 2001 From: AhmedAmraniAkdi Date: Tue, 21 May 2024 16:28:47 +0100 Subject: [PATCH 7/7] Skip test is alignment is 1 --- .../test_cl_ext_image_from_buffer.cpp | 68 ++++++++++--------- 1 file changed, 37 insertions(+), 31 deletions(-) diff --git a/test_conformance/images/kernel_read_write/test_cl_ext_image_from_buffer.cpp b/test_conformance/images/kernel_read_write/test_cl_ext_image_from_buffer.cpp index 7b4860dbcb..072acb9507 100644 --- a/test_conformance/images/kernel_read_write/test_cl_ext_image_from_buffer.cpp +++ b/test_conformance/images/kernel_read_write/test_cl_ext_image_from_buffer.cpp @@ -494,10 +494,11 @@ int image_from_buffer_alignment_negative(cl_device_id device, test_error(err, "Unable to create buffer"); /* Test Row pitch images */ - if (imageType == CL_MEM_OBJECT_IMAGE2D - || imageType == CL_MEM_OBJECT_IMAGE3D - || imageType == CL_MEM_OBJECT_IMAGE1D_ARRAY - || imageType == CL_MEM_OBJECT_IMAGE2D_ARRAY) + if ((imageType == CL_MEM_OBJECT_IMAGE2D + || imageType == CL_MEM_OBJECT_IMAGE3D + || imageType == CL_MEM_OBJECT_IMAGE1D_ARRAY + || imageType == CL_MEM_OBJECT_IMAGE2D_ARRAY) + && row_pitch_alignment != 1) { image_desc.buffer = buffer; image_desc.image_row_pitch = @@ -510,8 +511,9 @@ int image_from_buffer_alignment_negative(cl_device_id device, } /* Test Slice pitch images */ - if (imageType == CL_MEM_OBJECT_IMAGE3D - || imageType == CL_MEM_OBJECT_IMAGE2D_ARRAY) + if ((imageType == CL_MEM_OBJECT_IMAGE3D + || imageType == CL_MEM_OBJECT_IMAGE2D_ARRAY) + && slice_pitch_alignment != 1) { image_desc.buffer = buffer; image_desc.image_row_pitch = row_pitch; @@ -524,36 +526,40 @@ int image_from_buffer_alignment_negative(cl_device_id device, "Unexpected clCreateImage return"); } - /* Test buffer from host ptr to test base address alignment */ - const size_t aligned_buffer_size = - aligned_size(buffer_size, base_address_alignment); - /* Create buffer with host ptr and additional size for the wrong - * alignment */ - void* const host_ptr = - malloc(aligned_buffer_size + base_address_alignment); - void* non_aligned_host_ptr = - (void*)((char*)(aligned_ptr(host_ptr, - base_address_alignment)) - + 1); /* wrong alignment */ - - cl_mem buffer_host = clCreateBuffer( - context, CL_MEM_USE_HOST_PTR | CL_MEM_READ_WRITE, - buffer_size, non_aligned_host_ptr, &err); - test_error(err, "Unable to create buffer"); + if (base_address_alignment != 1) + { + /* Test buffer from host ptr to test base address alignment + */ + const size_t aligned_buffer_size = + aligned_size(buffer_size, base_address_alignment); + /* Create buffer with host ptr and additional size for the + * wrong alignment */ + void* const host_ptr = + malloc(aligned_buffer_size + base_address_alignment); + void* non_aligned_host_ptr = + (void*)((char*)(aligned_ptr(host_ptr, + base_address_alignment)) + + 1); /* wrong alignment */ + + cl_mem buffer_host = clCreateBuffer( + context, CL_MEM_USE_HOST_PTR | CL_MEM_READ_WRITE, + buffer_size, non_aligned_host_ptr, &err); + test_error(err, "Unable to create buffer"); - image_desc.buffer = buffer_host; + image_desc.buffer = buffer_host; - clCreateImage(context, flag, &format, &image_desc, nullptr, - &err); - test_failure_error(err, CL_INVALID_IMAGE_FORMAT_DESCRIPTOR, - "Unexpected clCreateImage return"); + clCreateImage(context, flag, &format, &image_desc, nullptr, + &err); + test_failure_error(err, CL_INVALID_IMAGE_FORMAT_DESCRIPTOR, + "Unexpected clCreateImage return"); - free(host_ptr); + free(host_ptr); - err = clReleaseMemObject(buffer); - test_error(err, "Unable to release buffer"); + err = clReleaseMemObject(buffer_host); + test_error(err, "Unable to release buffer"); + } - err = clReleaseMemObject(buffer_host); + err = clReleaseMemObject(buffer); test_error(err, "Unable to release buffer"); } }