diff --git a/test_conformance/profiling/CMakeLists.txt b/test_conformance/profiling/CMakeLists.txt index 036de05b..65772985 100644 --- a/test_conformance/profiling/CMakeLists.txt +++ b/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) diff --git a/test_conformance/profiling/main.cpp b/test_conformance/profiling/main.cpp index 04984728..012786cc 100644 --- a/test_conformance/profiling/main.cpp +++ b/test_conformance/profiling/main.cpp @@ -26,37 +26,38 @@ // the following variables (): 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 ); diff --git a/test_conformance/profiling/procs.h b/test_conformance/profiling/procs.h index cc7083d0..aa91a751 100644 --- a/test_conformance/profiling/procs.h +++ b/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__ diff --git a/test_conformance/profiling/profiling_timebase.cpp b/test_conformance/profiling/profiling_timebase.cpp new file mode 100644 index 00000000..1b12464c --- /dev/null +++ b/test_conformance/profiling/profiling_timebase.cpp @@ -0,0 +1,101 @@ +// +// 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"); + + err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, NULL, NULL, 1, &uEvent, + &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; +}