Skip to content

Commit

Permalink
Add cl_khr_kernel_clock tests (#1960)
Browse files Browse the repository at this point in the history
Adds cl_khr_kernel_clock test.

Also fixes failure in the compiler defines for extension compiler
subtest when cl_khr_kernel_clock is supported.
  • Loading branch information
AhmedAmraniAkdi authored May 21, 2024
1 parent 5ce18c3 commit cdf8d5e
Show file tree
Hide file tree
Showing 6 changed files with 269 additions and 1 deletion.
Original file line number Diff line number Diff line change
Expand Up @@ -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",
Expand Down Expand Up @@ -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",
Expand Down
1 change: 1 addition & 0 deletions test_conformance/extensions/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -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()
Original file line number Diff line number Diff line change
@@ -0,0 +1,8 @@
set(MODULE_NAME CL_KHR_KERNEL_CLOCK)

set(${MODULE_NAME}_SOURCES
main.cpp
kernel_clock.cpp
)

include(../../CMakeCommon.txt)
202 changes: 202 additions & 0 deletions test_conformance/extensions/cl_khr_kernel_clock/kernel_clock.cpp
Original file line number Diff line number Diff line change
@@ -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);
}
29 changes: 29 additions & 0 deletions test_conformance/extensions/cl_khr_kernel_clock/main.cpp
Original file line number Diff line number Diff line change
@@ -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);
}
27 changes: 27 additions & 0 deletions test_conformance/extensions/cl_khr_kernel_clock/procs.h
Original file line number Diff line number Diff line change
@@ -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*/

0 comments on commit cdf8d5e

Please sign in to comment.