Add cl_khr_kernel_clock tests (KhronosGroup#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_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",

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()

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)

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);
+}

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);
+}

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 <CL/cl.h>
+
+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*/