Queue profiling and Device timebase equality test

test_conformance/profiling/CMakeLists.txt

@@ -9,6 +9,7 @@ set(${MODULE_NAME}_SOURCES
  copy.cpp
  execute.cpp
  execute_multipass.cpp
+ profiling_timebase.cpp
 )
 
 include(../CMakeCommon.txt)

test_conformance/profiling/main.cpp

@@ -26,37 +26,38 @@
 // the following variables (<rdar://problem/11111245>):
 
 test_definition test_list[] = {
- ADD_TEST( read_array_int ),
- ADD_TEST( read_array_uint ),
- ADD_TEST( read_array_long ),
- ADD_TEST( read_array_ulong ),
- ADD_TEST( read_array_short ),
- ADD_TEST( read_array_ushort ),
- ADD_TEST( read_array_float ),
- ADD_TEST( read_array_char ),
- ADD_TEST( read_array_uchar ),
- ADD_TEST( read_array_struct ),
- ADD_TEST( write_array_int ),
- ADD_TEST( write_array_uint ),
- ADD_TEST( write_array_long ),
- ADD_TEST( write_array_ulong ),
- ADD_TEST( write_array_short ),
- ADD_TEST( write_array_ushort ),
- ADD_TEST( write_array_float ),
- ADD_TEST( write_array_char ),
- ADD_TEST( write_array_uchar ),
- ADD_TEST( write_array_struct ),
- ADD_TEST( read_image_float ),
- ADD_TEST( read_image_char ),
- ADD_TEST( read_image_uchar ),
- ADD_TEST( write_image_float ),
- ADD_TEST( write_image_char ),
- ADD_TEST( write_image_uchar ),
- ADD_TEST( copy_array ),
- ADD_TEST( copy_partial_array ),
- ADD_TEST( copy_image ),
- ADD_TEST( copy_array_to_image ),
- ADD_TEST( execute ),
+ ADD_TEST(read_array_int),
+ ADD_TEST(read_array_uint),
+ ADD_TEST(read_array_long),
+ ADD_TEST(read_array_ulong),
+ ADD_TEST(read_array_short),
+ ADD_TEST(read_array_ushort),
+ ADD_TEST(read_array_float),
+ ADD_TEST(read_array_char),
+ ADD_TEST(read_array_uchar),
+ ADD_TEST(read_array_struct),
+ ADD_TEST(write_array_int),
+ ADD_TEST(write_array_uint),
+ ADD_TEST(write_array_long),
+ ADD_TEST(write_array_ulong),
+ ADD_TEST(write_array_short),
+ ADD_TEST(write_array_ushort),
+ ADD_TEST(write_array_float),
+ ADD_TEST(write_array_char),
+ ADD_TEST(write_array_uchar),
+ ADD_TEST(write_array_struct),
+ ADD_TEST(read_image_float),
+ ADD_TEST(read_image_char),
+ ADD_TEST(read_image_uchar),
+ ADD_TEST(write_image_float),
+ ADD_TEST(write_image_char),
+ ADD_TEST(write_image_uchar),
+ ADD_TEST(copy_array),
+ ADD_TEST(copy_partial_array),
+ ADD_TEST(copy_image),
+ ADD_TEST(copy_array_to_image),
+ ADD_TEST(execute),
+ ADD_TEST_VERSION(profiling_timebase, Version(2, 1)),
 };
 
 const int test_num = ARRAY_SIZE( test_list );

test_conformance/profiling/procs.h

@@ -58,6 +58,8 @@ extern int test_copy_image( cl_device_id device, cl_context context, cl_c
 extern int test_copy_array_to_image( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements );
 extern int test_execute( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements );
 extern int test_parallel_kernels( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements );
+extern int test_profiling_timebase(cl_device_id device, cl_context context,
+ cl_command_queue queue, int num_elements);
 
 
 #endif // #ifndef __PROCS_H__

test_conformance/profiling/profiling_timebase.cpp

@@ -0,0 +1,102 @@
+//
+// Copyright (c) 2021 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"
+
+const char *kernelCode = "__kernel void kernel_empty(){}";
+
+int test_profiling_timebase(cl_device_id device, cl_context context,
+ cl_command_queue queue, int num_elements)
+{
+ Version version = get_device_cl_version(device);
+ cl_platform_id platform = getPlatformFromDevice(device);
+ cl_ulong timer_resolution = 0;
+ cl_int err =
+ clGetPlatformInfo(platform, CL_PLATFORM_HOST_TIMER_RESOLUTION,
+ sizeof(timer_resolution), &timer_resolution, NULL);
+ test_error(err, "Unable to query CL_PLATFORM_HOST_TIMER_RESOLUTION");
+
+ // If CL_PLATFORM_HOST_TIMER_RESOLUTION returns 0, clGetDeviceAndHostTimer
+ // is not a supported feature
+ if (timer_resolution == 0 && version >= Version(3, 0))
+ {
+ return TEST_SKIPPED_ITSELF;
+ }
+
+ cl_ulong hostTime;
+ clProgramWrapper program;
+ clKernelWrapper kernel;
+ clEventWrapper kEvent;
+
+ clEventWrapper uEvent = clCreateUserEvent(context, &err);
+ test_error(err, "Failed to create user event");
+
+ err = create_single_kernel_helper(context, &program, &kernel, 1,
+ &kernelCode, "kernel_empty");
+ test_error(err, "Failed to create kernel");
+
+ cl_ulong deviceTimeBeforeQueue;
+ err = clGetDeviceAndHostTimer(device, &deviceTimeBeforeQueue, &hostTime);
+ test_error(err, "Unable to get starting device time");
+
+ clEventWrapper userEvents = uEvent;
+ err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, NULL, NULL, 1,
+ &userEvents, &kEvent);
+ test_error(err, "clEnqueueNDRangeKernel failed");
+
+ cl_ulong deviceTimeAfterQueue;
+ err = clGetDeviceAndHostTimer(device, &deviceTimeAfterQueue, &hostTime);
+ test_error(err, "Unable to get queue device time");
+
+ err = clFlush(queue);
+ test_error(err, "clFlush failed");
+
+ err = clSetUserEventStatus(uEvent, CL_COMPLETE);
+ test_error(err, "Unable to complete user event");
+
+ err = clWaitForEvents(1, &kEvent);
+ test_error(err, "clWaitForEvents failed");
+
+ cl_ulong deviceTimeAfterCompletion;
+ err =
+ clGetDeviceAndHostTimer(device, &deviceTimeAfterCompletion, &hostTime);
+ test_error(err, "Unable to get finishing device time");
+
+ cl_ulong eventQueue, eventSubmit, eventStart, eventEnd;
+ err = clGetEventProfilingInfo(kEvent, CL_PROFILING_COMMAND_QUEUED,
+ sizeof(cl_ulong), &eventQueue, NULL);
+ test_error(err, "clGetEventProfilingInfo failed");
+ err = clGetEventProfilingInfo(kEvent, CL_PROFILING_COMMAND_SUBMIT,
+ sizeof(cl_ulong), &eventSubmit, NULL);
+ test_error(err, "clGetEventProfilingInfo failed");
+ err = clGetEventProfilingInfo(kEvent, CL_PROFILING_COMMAND_START,
+ sizeof(cl_ulong), &eventStart, NULL);
+ test_error(err, "clGetEventProfilingInfo failed");
+ err = clGetEventProfilingInfo(kEvent, CL_PROFILING_COMMAND_END,
+ sizeof(cl_ulong), &eventEnd, NULL);
+ test_error(err, "clGetEventProfilingInfo failed");
+
+ test_assert_error(deviceTimeBeforeQueue < eventQueue,
+ "Device timestamp was taken before kernel was queued");
+ test_assert_error(eventQueue < deviceTimeAfterQueue,
+ "Device timestamp was taken after kernel was queued");
+ test_assert_error(eventSubmit < deviceTimeAfterCompletion,
+ "Device timestamp was taken after kernel was submitted");
+ test_assert_error((eventStart < deviceTimeAfterCompletion)
+ && (eventEnd < deviceTimeAfterCompletion),
+ "Device timestamp was taken after kernel was executed");
+ return TEST_PASS;
+}