From cdf8d5e35ecc73e46278a36382c538cdb1ab2994 Mon Sep 17 00:00:00 2001 From: Ahmed <36049290+AhmedAmraniAkdi@users.noreply.github.com> Date: Tue, 21 May 2024 16:49:49 +0100 Subject: [PATCH] Add cl_khr_kernel_clock tests (#1960) Adds cl_khr_kernel_clock test. Also fixes failure in the compiler defines for extension compiler subtest when cl_khr_kernel_clock is supported. --- .../test_compiler_defines_for_extensions.cpp | 3 +- test_conformance/extensions/CMakeLists.txt | 1 + .../cl_khr_kernel_clock/CMakeLists.txt | 8 + .../cl_khr_kernel_clock/kernel_clock.cpp | 202 ++++++++++++++++++ .../extensions/cl_khr_kernel_clock/main.cpp | 29 +++ .../extensions/cl_khr_kernel_clock/procs.h | 27 +++ 6 files changed, 269 insertions(+), 1 deletion(-) 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/compiler/test_compiler_defines_for_extensions.cpp b/test_conformance/compiler/test_compiler_defines_for_extensions.cpp index 623c2dbfb7..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", @@ -94,7 +95,7 @@ const char *known_extensions[] = { }; 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/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..066ebb65d9 --- /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) 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..d8298b3d63 --- /dev/null +++ b/test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp @@ -0,0 +1,202 @@ +// 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" + +namespace { + +// write 1 to the output if the clock did not increase +static const char *kernel_sources[2] = { + 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 { + +public: + KernelClockTest(cl_device_id device, cl_context context, + cl_command_queue queue, + cl_device_kernel_clock_capabilities_khr capability) + : device(device), context(context), queue(queue), capability(capability) + {} + + 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); + 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_device_kernel_clock_capabilities_khr capability; +}; + +int MakeAndRunTest(cl_device_id device, cl_context context, + cl_command_queue queue, + cl_device_kernel_clock_capabilities_khr 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..a82564bc81 --- /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 + +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*/