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_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*/