ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-19 06:09:01 +00:00
Code Issues Packages Projects Releases Wiki Activity

Update test_external_semaphore.cpp with new kernels

Browse Source
This commit is contained in:
paulfradgley
2025-10-28 16:16:11 +00:00
committed by GitHub
parent b905f048d3
commit ae67dcff25
1 changed files with 100 additions and 10 deletions
Download Patch File Download Diff File Expand all files Collapse all files

110
test_conformance/extensions/cl_khr_external_semaphore/test_external_semaphore.cpp
View File

@@ -54,6 +54,41 @@
static const char *source = "__kernel void empty() {}";
const char* source_write_int = "__kernel void write_int(__global int* out, int val) { out[0] = val; }";
#define CREATE_KERNEL \
clProgramWrapper program_write_int; \
clKernelWrapper kernel_write_int; \
err = create_single_kernel_helper(context, &program_write_int, &kernel_write_int, 1, \
&source_write_int, "write_int"); \
size_t threads = 1; \
test_error(err, "Could not create kernel")
#define CREATE_BUFFER \
int int_val = 45; \
clMemWrapper buffer_write_int = clCreateBuffer(context, CL_MEM_READ_WRITE, \
sizeof(cl_int), nullptr, &err); \
test_error(err, "clCreateBuffer failed")
#define ENQUEUE_KERNEL(QUEUE, NUM_LIST, WAITLIST) \
err = clSetKernelArg(kernel_write_int, 0, sizeof(buffer_write_int), &buffer_write_int); \
test_error(err, "clSetKernelArg failed"); \
err = clSetKernelArg(kernel_write_int, 1, sizeof(int), &int_val); \
test_error(err, "clSetKernelArg failed"); \
err = clEnqueueNDRangeKernel(QUEUE, kernel_write_int, 1, nullptr, \
&threads, nullptr, NUM_LIST, WAITLIST, nullptr); \
test_error(err, "clEnqueueNDRangeKernel failed")
#define ENQUEUE_KERNEL_WITH_EVENT(QUEUE, NUM_LIST, WAITLIST, EVENT) \
err = clSetKernelArg(kernel_write_int, 0, sizeof(buffer_write_int), &buffer_write_int); \
test_error(err, "clSetKernelArg failed"); \
err = clSetKernelArg(kernel_write_int, 1, sizeof(int), &int_val); \
test_error(err, "clSetKernelArg failed"); \
clEventWrapper EVENT; \
err = clEnqueueNDRangeKernel(QUEUE, kernel_write_int, 1, nullptr, \
&threads, nullptr, NUM_LIST, WAITLIST, &EVENT); \
test_error(err, "clEnqueueNDRangeKernel failed")
static void log_info_semaphore_type(
VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType)
{
@@ -249,6 +284,9 @@ REGISTER_TEST_VERSION(external_semaphores_cross_context, Version(1, 2))
clCreateCommandQueue(context2, device, 0, &err);
test_error(err, "Could not create command queue");
CREATE_KERNEL;
CREATE_BUFFER;
if (import_export_handle_types.empty())
{
log_info("Could not find a handle type that supports both import and "
@@ -268,13 +306,15 @@ REGISTER_TEST_VERSION(external_semaphores_cross_context, Version(1, 2))
(cl_semaphore_properties_khr)0
};
ENQUEUE_KERNEL_WITH_EVENT(queue1, 0, nullptr, write_int_event);
// Signal semaphore on context1
cl_semaphore_khr exportable_semaphore =
clCreateSemaphoreWithPropertiesKHR(context, export_props, &err);
test_error(err, "Failed to create exportable semaphore");
err = clEnqueueSignalSemaphoresKHR(queue1, 1, &exportable_semaphore,
nullptr, 0, nullptr, nullptr);
nullptr, 1, &write_int_event, nullptr);
test_error(err, "Failed to signal semaphore on context1");
cl_semaphore_properties_khr handle =
@@ -296,10 +336,13 @@ REGISTER_TEST_VERSION(external_semaphores_cross_context, Version(1, 2))
clCreateSemaphoreWithPropertiesKHR(context2, import_props, &err);
test_error(err, "Failed to import semaphore into context2 semaphore");
clEventWrapper wait_event;
err = clEnqueueWaitSemaphoresKHR(queue2, 1, &imported_semaphore,
nullptr, 0, nullptr, nullptr);
nullptr, 0, nullptr, &wait_event);
test_error(err, "Failed to signal semaphore on context1");
ENQUEUE_KERNEL(queue2, 1, &wait_event);
err = clFlush(queue1);
test_error(err, "Failed to flush queue1");
@@ -346,6 +389,9 @@ REGISTER_TEST_VERSION(external_semaphores_simple_1, Version(1, 2))
test_fail("No external semaphore handle types found\n");
}
CREATE_KERNEL;
CREATE_BUFFER;
for (VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType :
vkExternalSemaphoreHandleTypeList)
{
@@ -364,10 +410,12 @@ REGISTER_TEST_VERSION(external_semaphores_simple_1, Version(1, 2))
context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
test_error(err, "Could not create command queue");
ENQUEUE_KERNEL_WITH_EVENT(queue, 0, nullptr, write_int_event);
// Signal semaphore
clEventWrapper signal_event;
err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext.getCLSemaphore(),
nullptr, 0, nullptr, &signal_event);
nullptr, 1, &write_int_event, &signal_event);
test_error(err, "Could not signal semaphore");
// Wait semaphore
@@ -376,6 +424,8 @@ REGISTER_TEST_VERSION(external_semaphores_simple_1, Version(1, 2))
nullptr, 0, nullptr, &wait_event);
test_error(err, "Could not wait semaphore");
ENQUEUE_KERNEL(queue, 1, &wait_event);
// Finish
err = clFinish(queue);
test_error(err, "Could not finish queue");
@@ -415,6 +465,9 @@ REGISTER_TEST_VERSION(external_semaphores_reuse, Version(1, 2))
test_fail("No external semaphore handle types found\n");
}
CREATE_KERNEL;
CREATE_BUFFER;
for (VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType :
vkExternalSemaphoreHandleTypeList)
{
@@ -486,6 +539,8 @@ REGISTER_TEST_VERSION(external_semaphores_reuse, Version(1, 2))
&wait_events[loop - 1]);
test_error(err, "Could not wait semaphore");
ENQUEUE_KERNEL(queue, 1, &wait_events[loop - 1]);
// Finish
err = clFinish(queue);
test_error(err, "Could not finish queue");
@@ -533,6 +588,9 @@ static int external_semaphore_cross_queue_helper(cl_device_id device,
test_fail("No external semaphore handle types found\n");
}
CREATE_KERNEL;
CREATE_BUFFER;
for (VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType :
vkExternalSemaphoreHandleTypeList)
{
@@ -545,11 +603,13 @@ static int external_semaphore_cross_queue_helper(cl_device_id device,
cl_int err = CL_SUCCESS;
ENQUEUE_KERNEL_WITH_EVENT(queue_1, 0, nullptr, write_int_event);
// Signal semaphore on queue_1
clEventWrapper signal_event;
err =
clEnqueueSignalSemaphoresKHR(queue_1, 1, &sema_ext.getCLSemaphore(),
nullptr, 0, nullptr, &signal_event);
nullptr, 1, &write_int_event, &signal_event);
test_error(err, "Could not signal semaphore");
// Wait semaphore on queue_2
@@ -558,6 +618,8 @@ static int external_semaphore_cross_queue_helper(cl_device_id device,
nullptr, 0, nullptr, &wait_event);
test_error(err, "Could not wait semaphore");
ENQUEUE_KERNEL(queue_2, 1, &wait_event);
// Finish queue_1 and queue_2
err = clFinish(queue_1);
test_error(err, "Could not finish queue");
@@ -645,6 +707,9 @@ REGISTER_TEST_VERSION(external_semaphores_cross_queues_io2, Version(1, 2))
test_fail("No external semaphore handle types found\n");
}
CREATE_KERNEL;
CREATE_BUFFER;
for (VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType :
vkExternalSemaphoreHandleTypeList)
{
@@ -666,10 +731,12 @@ REGISTER_TEST_VERSION(external_semaphores_cross_queues_io2, Version(1, 2))
clCreateCommandQueue(context2, device, 0, &err);
test_error(err, "Could not create command queue");
ENQUEUE_KERNEL_WITH_EVENT(queue1, 0, nullptr, write_int_event);
// Signal semaphore 1
clEventWrapper signal_1_event;
err = clEnqueueSignalSemaphoresKHR(
queue1, 1, &sema_ext_1.getCLSemaphore(), nullptr, 0, nullptr,
queue1, 1, &sema_ext_1.getCLSemaphore(), nullptr, 1, &write_int_event,
&signal_1_event);
test_error(err, "Could not signal semaphore");
@@ -680,10 +747,13 @@ REGISTER_TEST_VERSION(external_semaphores_cross_queues_io2, Version(1, 2))
nullptr, 0, nullptr, &wait_1_event);
test_error(err, "Could not wait semaphore");
ENQUEUE_KERNEL(queue1, 1, &wait_1_event);
ENQUEUE_KERNEL(queue2, 0, nullptr, write_int_event_2);
// Signal semaphore 2
clEventWrapper signal_2_event;
err = clEnqueueSignalSemaphoresKHR(
queue2, 1, &sema_ext_2.getCLSemaphore(), nullptr, 0, nullptr,
queue2, 1, &sema_ext_2.getCLSemaphore(), nullptr, 1, &write_int_event_2,
&signal_2_event);
test_error(err, "Could not signal semaphore");
@@ -694,6 +764,8 @@ REGISTER_TEST_VERSION(external_semaphores_cross_queues_io2, Version(1, 2))
nullptr, 0, nullptr, &wait_2_event);
test_error(err, "Could not wait semaphore");
ENQUEUE_KERNEL(queue2, 1, &wait_2_event);
// Finish
err = clFinish(queue1);
test_error(err, "Could not finish queue");
@@ -738,6 +810,9 @@ REGISTER_TEST_VERSION(external_semaphores_multi_signal, Version(1, 2))
test_fail("No external semaphore handle types found\n");
}
CREATE_KERNEL;
CREATE_BUFFER;
for (VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType :
vkExternalSemaphoreHandleTypeList)
{
@@ -760,12 +835,14 @@ REGISTER_TEST_VERSION(external_semaphores_multi_signal, Version(1, 2))
context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
test_error(err, "Could not create command queue");
ENQUEUE_KERNEL_WITH_EVENT(queue, 0, nullptr, write_int_event);
// Signal semaphore 1 and 2
clEventWrapper signal_event;
cl_semaphore_khr sema_list[] = { sema_ext_1.getCLSemaphore(),
sema_ext_2.getCLSemaphore() };
err = clEnqueueSignalSemaphoresKHR(queue, 2, sema_list, nullptr, 0,
nullptr, &signal_event);
err = clEnqueueSignalSemaphoresKHR(queue, 2, sema_list, nullptr, 1,
&write_int_event, &signal_event);
test_error(err, "Could not signal semaphore");
// Wait semaphore 1
@@ -774,11 +851,15 @@ REGISTER_TEST_VERSION(external_semaphores_multi_signal, Version(1, 2))
nullptr, 0, nullptr, &wait_1_event);
test_error(err, "Could not wait semaphore");
ENQUEUE_KERNEL(queue, 1, &wait_1_event);
// Wait semaphore 2
clEventWrapper wait_2_event;
err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_ext_2.getCLSemaphore(),
nullptr, 0, nullptr, &wait_2_event);
test_error(err, "Could not wait semaphore");
ENQUEUE_KERNEL(queue, 1, &wait_2_event);
// Finish
err = clFinish(queue);
@@ -820,6 +901,9 @@ REGISTER_TEST_VERSION(external_semaphores_multi_wait, Version(1, 2))
test_fail("No external semaphore handle types found\n");
}
CREATE_KERNEL;
CREATE_BUFFER;
for (VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType :
vkExternalSemaphoreHandleTypeList)
{
@@ -842,18 +926,22 @@ REGISTER_TEST_VERSION(external_semaphores_multi_wait, Version(1, 2))
context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
test_error(err, "Could not create command queue");
ENQUEUE_KERNEL_WITH_EVENT(queue, 0, nullptr, write_int_event_1);
// Signal semaphore 1
clEventWrapper signal_1_event;
err =
clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext_1.getCLSemaphore(),
nullptr, 0, nullptr, &signal_1_event);
nullptr, 1, write_int_event_1, &signal_1_event);
test_error(err, "Could not signal semaphore");
ENQUEUE_KERNEL_WITH_EVENT(queue, 0, nullptr, write_int_event_2);
// Signal semaphore 2
clEventWrapper signal_2_event;
err =
clEnqueueSignalSemaphoresKHR(queue, 1, &sema_ext_2.getCLSemaphore(),
nullptr, 0, nullptr, &signal_2_event);
nullptr, 1, write_int_event_2, &signal_2_event);
test_error(err, "Could not signal semaphore");
// Wait semaphore 1 and 2
@@ -864,6 +952,8 @@ REGISTER_TEST_VERSION(external_semaphores_multi_wait, Version(1, 2))
nullptr, &wait_event);
test_error(err, "Could not wait semaphore");
ENQUEUE_KERNEL(queue, 1, &wait_event);
// Finish
err = clFinish(queue);
test_error(err, "Could not finish queue");