From ae67dcff250055d0a1fec5e12a7232000fa37536 Mon Sep 17 00:00:00 2001 From: paulfradgley <39525348+paulfradgley@users.noreply.github.com> Date: Tue, 28 Oct 2025 16:16:11 +0000 Subject: [PATCH] Update test_external_semaphore.cpp with new kernels --- .../test_external_semaphore.cpp | 110 ++++++++++++++++-- 1 file changed, 100 insertions(+), 10 deletions(-) diff --git a/test_conformance/extensions/cl_khr_external_semaphore/test_external_semaphore.cpp b/test_conformance/extensions/cl_khr_external_semaphore/test_external_semaphore.cpp index 198bf046..87ef177e 100644 --- a/test_conformance/extensions/cl_khr_external_semaphore/test_external_semaphore.cpp +++ b/test_conformance/extensions/cl_khr_external_semaphore/test_external_semaphore.cpp @@ -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");