1 //
2 // Copyright (c) 2023 The Khronos Group Inc.
3 //
4 // Licensed under the Apache License, Version 2.0 (the "License");
5 // you may not use this file except in compliance with the License.
6 // You may obtain a copy of the License at
7 //
8 // http://www.apache.org/licenses/LICENSE-2.0
9 //
10 // Unless required by applicable law or agreed to in writing, software
11 // distributed under the License is distributed on an "AS IS" BASIS,
12 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 // See the License for the specific language governing permissions and
14 // limitations under the License.
15 //
16
17
18 #include "harness/typeWrappers.h"
19 #include "harness/extensionHelpers.h"
20 #include "harness/errorHelpers.h"
21 #include <system_error>
22 #include <thread>
23 #include <chrono>
24
25 #define FLUSH_DELAY_S 5
26
27 #define SEMAPHORE_PARAM_TEST(param_name, param_type, expected) \
28 do \
29 { \
30 param_type value; \
31 size_t size; \
32 cl_int error = clGetSemaphoreInfoKHR(sema, param_name, sizeof(value), \
33 &value, &size); \
34 test_error(error, "Unable to get " #param_name " from semaphore"); \
35 if (value != expected) \
36 { \
37 test_fail("ERROR: Parameter %s did not validate! (expected %d, " \
38 "got %d)\n", \
39 #param_name, expected, value); \
40 } \
41 if (size != sizeof(value)) \
42 { \
43 test_fail( \
44 "ERROR: Returned size of parameter %s does not validate! " \
45 "(expected %d, got %d)\n", \
46 #param_name, (int)sizeof(value), (int)size); \
47 } \
48 } while (false)
49
50 #define SEMAPHORE_PARAM_TEST_ARRAY(param_name, param_type, num_params, \
51 expected) \
52 do \
53 { \
54 param_type value[num_params]; \
55 size_t size; \
56 cl_int error = clGetSemaphoreInfoKHR(sema, param_name, sizeof(value), \
57 &value, &size); \
58 test_error(error, "Unable to get " #param_name " from semaphore"); \
59 if (size != sizeof(value)) \
60 { \
61 test_fail( \
62 "ERROR: Returned size of parameter %s does not validate! " \
63 "(expected %d, got %d)\n", \
64 #param_name, (int)sizeof(value), (int)size); \
65 } \
66 if (memcmp(value, expected, size) != 0) \
67 { \
68 test_fail("ERROR: Parameter %s did not validate!\n", #param_name); \
69 } \
70 } while (false)
71
72 static const char* source = "__kernel void empty() {}";
73
74 // Helper function that signals and waits on semaphore across two different
75 // queues.
semaphore_cross_queue_helper(cl_device_id deviceID,cl_context context,cl_command_queue queue_1,cl_command_queue queue_2)76 static int semaphore_cross_queue_helper(cl_device_id deviceID,
77 cl_context context,
78 cl_command_queue queue_1,
79 cl_command_queue queue_2)
80 {
81 cl_int err;
82
83 if (!is_extension_available(deviceID, "cl_khr_semaphore"))
84 {
85 log_info("cl_khr_semaphore is not supported on this platoform. "
86 "Skipping test.\n");
87 return TEST_SKIPPED_ITSELF;
88 }
89
90 // Obtain pointers to semaphore's API
91 GET_PFN(deviceID, clCreateSemaphoreWithPropertiesKHR);
92 GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR);
93 GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR);
94 GET_PFN(deviceID, clReleaseSemaphoreKHR);
95
96 // Create semaphore
97 cl_semaphore_properties_khr sema_props[] = {
98 static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_KHR),
99 static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_BINARY_KHR),
100 0
101 };
102 cl_semaphore_khr sema =
103 clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err);
104 test_error(err, "Could not create semaphore");
105
106 // Signal semaphore on queue_1
107 clEventWrapper signal_event;
108 err = clEnqueueSignalSemaphoresKHR(queue_1, 1, &sema, nullptr, 0, nullptr,
109 &signal_event);
110 test_error(err, "Could not signal semaphore");
111
112 // Wait semaphore on queue_2
113 clEventWrapper wait_event;
114 err = clEnqueueWaitSemaphoresKHR(queue_2, 1, &sema, nullptr, 0, nullptr,
115 &wait_event);
116 test_error(err, "Could not wait semaphore");
117
118 // Finish queue_1 and queue_2
119 err = clFinish(queue_1);
120 test_error(err, "Could not finish queue");
121
122 err = clFinish(queue_2);
123 test_error(err, "Could not finish queue");
124
125 // Ensure all events are completed
126 test_assert_event_complete(signal_event);
127 test_assert_event_complete(wait_event);
128
129 // Release semaphore
130 err = clReleaseSemaphoreKHR(sema);
131 test_error(err, "Could not release semaphore");
132
133 return TEST_PASS;
134 }
135
136 // Confirm that a signal followed by a wait will complete successfully
test_semaphores_simple_1(cl_device_id deviceID,cl_context context,cl_command_queue defaultQueue,int num_elements)137 int test_semaphores_simple_1(cl_device_id deviceID, cl_context context,
138 cl_command_queue defaultQueue, int num_elements)
139 {
140 cl_int err;
141
142 if (!is_extension_available(deviceID, "cl_khr_semaphore"))
143 {
144 log_info("cl_khr_semaphore is not supported on this platoform. "
145 "Skipping test.\n");
146 return TEST_SKIPPED_ITSELF;
147 }
148
149 // Obtain pointers to semaphore's API
150 GET_PFN(deviceID, clCreateSemaphoreWithPropertiesKHR);
151 GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR);
152 GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR);
153 GET_PFN(deviceID, clReleaseSemaphoreKHR);
154
155 // Create ooo queue
156 clCommandQueueWrapper queue = clCreateCommandQueue(
157 context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
158 test_error(err, "Could not create command queue");
159
160 // Create semaphore
161 cl_semaphore_properties_khr sema_props[] = {
162 static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_KHR),
163 static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_BINARY_KHR),
164 0
165 };
166 cl_semaphore_khr sema =
167 clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err);
168 test_error(err, "Could not create semaphore");
169
170 // Signal semaphore
171 clEventWrapper signal_event;
172 err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema, nullptr, 0, nullptr,
173 &signal_event);
174 test_error(err, "Could not signal semaphore");
175
176 // Wait semaphore
177 clEventWrapper wait_event;
178 err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema, nullptr, 0, nullptr,
179 &wait_event);
180 test_error(err, "Could not wait semaphore");
181
182 // Finish
183 err = clFinish(queue);
184 test_error(err, "Could not finish queue");
185
186 // Ensure all events are completed
187 test_assert_event_complete(signal_event);
188 test_assert_event_complete(wait_event);
189
190 // Release semaphore
191 err = clReleaseSemaphoreKHR(sema);
192 test_error(err, "Could not release semaphore");
193
194 return TEST_PASS;
195 }
196
197 // Confirm that signal a semaphore with no event dependencies will not result
198 // in an implicit dependency on everything previously submitted
test_semaphores_simple_2(cl_device_id deviceID,cl_context context,cl_command_queue defaultQueue,int num_elements)199 int test_semaphores_simple_2(cl_device_id deviceID, cl_context context,
200 cl_command_queue defaultQueue, int num_elements)
201 {
202 cl_int err;
203
204 if (!is_extension_available(deviceID, "cl_khr_semaphore"))
205 {
206 log_info("cl_khr_semaphore is not supported on this platoform. "
207 "Skipping test.\n");
208 return TEST_SKIPPED_ITSELF;
209 }
210
211 // Obtain pointers to semaphore's API
212 GET_PFN(deviceID, clCreateSemaphoreWithPropertiesKHR);
213 GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR);
214 GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR);
215 GET_PFN(deviceID, clReleaseSemaphoreKHR);
216
217 // Create ooo queue
218 clCommandQueueWrapper queue = clCreateCommandQueue(
219 context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
220 test_error(err, "Could not create command queue");
221
222 // Create semaphore
223 cl_semaphore_properties_khr sema_props[] = {
224 static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_KHR),
225 static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_BINARY_KHR),
226 0
227 };
228 cl_semaphore_khr sema =
229 clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err);
230 test_error(err, "Could not create semaphore");
231
232 // Create user event
233 clEventWrapper user_event = clCreateUserEvent(context, &err);
234 test_error(err, "Could not create user event");
235
236 // Create Kernel
237 clProgramWrapper program;
238 clKernelWrapper kernel;
239 err = create_single_kernel_helper(context, &program, &kernel, 1, &source,
240 "empty");
241 test_error(err, "Could not create kernel");
242
243 // Enqueue task_1 (dependency on user_event)
244 clEventWrapper task_1_event;
245 err = clEnqueueTask(queue, kernel, 1, &user_event, &task_1_event);
246 test_error(err, "Could not enqueue task 1");
247
248 // Signal semaphore
249 clEventWrapper signal_event;
250 err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema, nullptr, 0, nullptr,
251 &signal_event);
252 test_error(err, "Could not signal semaphore");
253
254 // Wait semaphore
255 clEventWrapper wait_event;
256 err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema, nullptr, 0, nullptr,
257 &wait_event);
258 test_error(err, "Could not wait semaphore");
259
260 // Flush and delay
261 err = clFlush(queue);
262 test_error(err, "Could not flush queue");
263 std::this_thread::sleep_for(std::chrono::seconds(FLUSH_DELAY_S));
264
265 // Ensure all events are completed except for task_1
266 test_assert_event_inprogress(task_1_event);
267 test_assert_event_complete(signal_event);
268 test_assert_event_complete(wait_event);
269
270 // Complete user_event
271 err = clSetUserEventStatus(user_event, CL_COMPLETE);
272 test_error(err, "Could not set user event to CL_COMPLETE");
273
274 // Finish
275 err = clFinish(queue);
276 test_error(err, "Could not finish queue");
277
278 // Ensure all events are completed
279 test_assert_event_complete(task_1_event);
280 test_assert_event_complete(signal_event);
281 test_assert_event_complete(wait_event);
282
283 // Release semaphore
284 err = clReleaseSemaphoreKHR(sema);
285 test_error(err, "Could not release semaphore");
286
287 return TEST_PASS;
288 }
289
290 // Confirm that a semaphore can be reused multiple times
test_semaphores_reuse(cl_device_id deviceID,cl_context context,cl_command_queue defaultQueue,int num_elements)291 int test_semaphores_reuse(cl_device_id deviceID, cl_context context,
292 cl_command_queue defaultQueue, int num_elements)
293 {
294 cl_int err;
295
296 if (!is_extension_available(deviceID, "cl_khr_semaphore"))
297 {
298 log_info("cl_khr_semaphore is not supported on this platoform. "
299 "Skipping test.\n");
300 return TEST_SKIPPED_ITSELF;
301 }
302
303 // Obtain pointers to semaphore's API
304 GET_PFN(deviceID, clCreateSemaphoreWithPropertiesKHR);
305 GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR);
306 GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR);
307 GET_PFN(deviceID, clReleaseSemaphoreKHR);
308
309 // Create ooo queue
310 clCommandQueueWrapper queue = clCreateCommandQueue(
311 context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
312 test_error(err, "Could not create command queue");
313
314 // Create semaphore
315 cl_semaphore_properties_khr sema_props[] = {
316 static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_KHR),
317 static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_BINARY_KHR),
318 0
319 };
320 cl_semaphore_khr sema =
321 clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err);
322 test_error(err, "Could not create semaphore");
323
324 // Create Kernel
325 clProgramWrapper program;
326 clKernelWrapper kernel;
327 err = create_single_kernel_helper(context, &program, &kernel, 1, &source,
328 "empty");
329 test_error(err, "Could not create kernel");
330
331 constexpr size_t loop_count = 10;
332 clEventWrapper signal_events[loop_count];
333 clEventWrapper wait_events[loop_count];
334 clEventWrapper task_events[loop_count];
335
336 // Enqueue task_1
337 err = clEnqueueTask(queue, kernel, 0, nullptr, &task_events[0]);
338 test_error(err, "Unable to enqueue task_1");
339
340 // Signal semaphore (dependency on task_1)
341 err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema, nullptr, 1,
342 &task_events[0], &signal_events[0]);
343 test_error(err, "Could not signal semaphore");
344
345 // In a loop
346 size_t loop;
347 for (loop = 1; loop < loop_count; ++loop)
348 {
349 // Wait semaphore
350 err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema, nullptr, 0, nullptr,
351 &wait_events[loop - 1]);
352 test_error(err, "Could not wait semaphore");
353
354 // Enqueue task_loop (dependency on wait)
355 err = clEnqueueTask(queue, kernel, 1, &wait_events[loop - 1],
356 &task_events[loop]);
357 test_error(err, "Unable to enqueue task_loop");
358
359 // Wait for the "wait semaphore" to complete
360 err = clWaitForEvents(1, &wait_events[loop - 1]);
361 test_error(err, "Unable to wait for wait semaphore to complete");
362
363 // Signal semaphore (dependency on task_loop)
364 err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema, nullptr, 1,
365 &task_events[loop],
366 &signal_events[loop]);
367 test_error(err, "Could not signal semaphore");
368 }
369
370 // Wait semaphore
371 err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema, nullptr, 0, nullptr,
372 &wait_events[loop - 1]);
373 test_error(err, "Could not wait semaphore");
374
375 // Finish
376 err = clFinish(queue);
377 test_error(err, "Could not finish queue");
378
379 // Ensure all events are completed
380 for (loop = 0; loop < loop_count; ++loop)
381 {
382 test_assert_event_complete(wait_events[loop]);
383 test_assert_event_complete(signal_events[loop]);
384 test_assert_event_complete(task_events[loop]);
385 }
386
387 // Release semaphore
388 err = clReleaseSemaphoreKHR(sema);
389 test_error(err, "Could not release semaphore");
390
391 return TEST_PASS;
392 }
393
394 // Confirm that a semaphore works across different ooo queues
test_semaphores_cross_queues_ooo(cl_device_id deviceID,cl_context context,cl_command_queue defaultQueue,int num_elements)395 int test_semaphores_cross_queues_ooo(cl_device_id deviceID, cl_context context,
396 cl_command_queue defaultQueue,
397 int num_elements)
398 {
399 cl_int err;
400
401 // Create ooo queues
402 clCommandQueueWrapper queue_1 = clCreateCommandQueue(
403 context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
404 test_error(err, "Could not create command queue");
405
406 clCommandQueueWrapper queue_2 = clCreateCommandQueue(
407 context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
408 test_error(err, "Could not create command queue");
409
410 return semaphore_cross_queue_helper(deviceID, context, queue_1, queue_2);
411 }
412
413 // Confirm that a semaphore works across different in-order queues
test_semaphores_cross_queues_io(cl_device_id deviceID,cl_context context,cl_command_queue defaultQueue,int num_elements)414 int test_semaphores_cross_queues_io(cl_device_id deviceID, cl_context context,
415 cl_command_queue defaultQueue,
416 int num_elements)
417 {
418 cl_int err;
419
420 // Create in-order queues
421 clCommandQueueWrapper queue_1 =
422 clCreateCommandQueue(context, deviceID, 0, &err);
423 test_error(err, "Could not create command queue");
424
425 clCommandQueueWrapper queue_2 =
426 clCreateCommandQueue(context, deviceID, 0, &err);
427 test_error(err, "Could not create command queue");
428
429 return semaphore_cross_queue_helper(deviceID, context, queue_1, queue_2);
430 }
431
432 // Confirm that we can signal multiple semaphores with one command
test_semaphores_multi_signal(cl_device_id deviceID,cl_context context,cl_command_queue defaultQueue,int num_elements)433 int test_semaphores_multi_signal(cl_device_id deviceID, cl_context context,
434 cl_command_queue defaultQueue,
435 int num_elements)
436 {
437 cl_int err;
438
439 if (!is_extension_available(deviceID, "cl_khr_semaphore"))
440 {
441 log_info("cl_khr_semaphore is not supported on this platoform. "
442 "Skipping test.\n");
443 return TEST_SKIPPED_ITSELF;
444 }
445
446 // Obtain pointers to semaphore's API
447 GET_PFN(deviceID, clCreateSemaphoreWithPropertiesKHR);
448 GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR);
449 GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR);
450 GET_PFN(deviceID, clReleaseSemaphoreKHR);
451
452 // Create ooo queue
453 clCommandQueueWrapper queue = clCreateCommandQueue(
454 context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
455 test_error(err, "Could not create command queue");
456
457 // Create semaphore
458 cl_semaphore_properties_khr sema_props[] = {
459 static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_KHR),
460 static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_BINARY_KHR),
461 0
462 };
463 cl_semaphore_khr sema_1 =
464 clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err);
465 test_error(err, "Could not create semaphore");
466
467 cl_semaphore_khr sema_2 =
468 clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err);
469 test_error(err, "Could not create semaphore");
470
471 // Signal semaphore 1 and 2
472 clEventWrapper signal_event;
473 cl_semaphore_khr sema_list[] = { sema_1, sema_2 };
474 err = clEnqueueSignalSemaphoresKHR(queue, 2, sema_list, nullptr, 0, nullptr,
475 &signal_event);
476 test_error(err, "Could not signal semaphore");
477
478 // Wait semaphore 1
479 clEventWrapper wait_1_event;
480 err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_1, nullptr, 0, nullptr,
481 &wait_1_event);
482 test_error(err, "Could not wait semaphore");
483
484 // Wait semaphore 2
485 clEventWrapper wait_2_event;
486 err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_2, nullptr, 0, nullptr,
487 &wait_2_event);
488 test_error(err, "Could not wait semaphore");
489
490 // Finish
491 err = clFinish(queue);
492 test_error(err, "Could not finish queue");
493
494 // Ensure all events are completed
495 test_assert_event_complete(signal_event);
496 test_assert_event_complete(wait_1_event);
497 test_assert_event_complete(wait_2_event);
498
499 // Release semaphores
500 err = clReleaseSemaphoreKHR(sema_1);
501 test_error(err, "Could not release semaphore");
502
503 err = clReleaseSemaphoreKHR(sema_2);
504 test_error(err, "Could not release semaphore");
505
506 return TEST_PASS;
507 }
508
509 // Confirm that we can wait for multiple semaphores with one command
test_semaphores_multi_wait(cl_device_id deviceID,cl_context context,cl_command_queue defaultQueue,int num_elements)510 int test_semaphores_multi_wait(cl_device_id deviceID, cl_context context,
511 cl_command_queue defaultQueue, int num_elements)
512 {
513 cl_int err;
514
515 if (!is_extension_available(deviceID, "cl_khr_semaphore"))
516 {
517 log_info("cl_khr_semaphore is not supported on this platoform. "
518 "Skipping test.\n");
519 return TEST_SKIPPED_ITSELF;
520 }
521
522 // Obtain pointers to semaphore's API
523 GET_PFN(deviceID, clCreateSemaphoreWithPropertiesKHR);
524 GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR);
525 GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR);
526 GET_PFN(deviceID, clReleaseSemaphoreKHR);
527
528 // Create ooo queue
529 clCommandQueueWrapper queue = clCreateCommandQueue(
530 context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
531 test_error(err, "Could not create command queue");
532
533 // Create semaphores
534 cl_semaphore_properties_khr sema_props[] = {
535 static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_KHR),
536 static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_BINARY_KHR),
537 0
538 };
539 cl_semaphore_khr sema_1 =
540 clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err);
541 test_error(err, "Could not create semaphore");
542
543 cl_semaphore_khr sema_2 =
544 clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err);
545 test_error(err, "Could not create semaphore");
546
547 // Signal semaphore 1
548 clEventWrapper signal_1_event;
549 err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema_1, nullptr, 0, nullptr,
550 &signal_1_event);
551 test_error(err, "Could not signal semaphore");
552
553 // Signal semaphore 2
554 clEventWrapper signal_2_event;
555 err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema_2, nullptr, 0, nullptr,
556 &signal_2_event);
557 test_error(err, "Could not signal semaphore");
558
559 // Wait semaphore 1 and 2
560 clEventWrapper wait_event;
561 cl_semaphore_khr sema_list[] = { sema_1, sema_2 };
562 err = clEnqueueWaitSemaphoresKHR(queue, 2, sema_list, nullptr, 0, nullptr,
563 &wait_event);
564 test_error(err, "Could not wait semaphore");
565
566 // Finish
567 err = clFinish(queue);
568 test_error(err, "Could not finish queue");
569
570 // Ensure all events are completed
571 test_assert_event_complete(signal_1_event);
572 test_assert_event_complete(signal_2_event);
573 test_assert_event_complete(wait_event);
574
575 // Release semaphores
576 err = clReleaseSemaphoreKHR(sema_1);
577 test_error(err, "Could not release semaphore");
578
579 err = clReleaseSemaphoreKHR(sema_2);
580 test_error(err, "Could not release semaphore");
581
582 return TEST_PASS;
583 }
584
585 // Confirm the semaphores can be successfully queried
test_semaphores_queries(cl_device_id deviceID,cl_context context,cl_command_queue defaultQueue,int num_elements)586 int test_semaphores_queries(cl_device_id deviceID, cl_context context,
587 cl_command_queue defaultQueue, int num_elements)
588 {
589 cl_int err;
590
591 if (!is_extension_available(deviceID, "cl_khr_semaphore"))
592 {
593 log_info("cl_khr_semaphore is not supported on this platoform. "
594 "Skipping test.\n");
595 return TEST_SKIPPED_ITSELF;
596 }
597
598 // Obtain pointers to semaphore's API
599 GET_PFN(deviceID, clCreateSemaphoreWithPropertiesKHR);
600 GET_PFN(deviceID, clGetSemaphoreInfoKHR);
601 GET_PFN(deviceID, clRetainSemaphoreKHR);
602 GET_PFN(deviceID, clReleaseSemaphoreKHR);
603
604 // Create binary semaphore
605 cl_semaphore_properties_khr sema_props[] = {
606 static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_KHR),
607 static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_BINARY_KHR),
608 0
609 };
610 cl_semaphore_khr sema =
611 clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err);
612 test_error(err, "Could not create semaphore");
613
614 // Confirm that querying CL_SEMAPHORE_TYPE_KHR returns
615 // CL_SEMAPHORE_TYPE_BINARY_KHR
616 SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_TYPE_KHR, cl_semaphore_type_khr,
617 CL_SEMAPHORE_TYPE_BINARY_KHR);
618
619 // Confirm that querying CL_SEMAPHORE_CONTEXT_KHR returns the right context
620 SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_CONTEXT_KHR, cl_context, context);
621
622 // Confirm that querying CL_SEMAPHORE_REFERENCE_COUNT_KHR returns the right
623 // value
624 SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_REFERENCE_COUNT_KHR, cl_uint, 1);
625
626 err = clRetainSemaphoreKHR(sema);
627 test_error(err, "Could not retain semaphore");
628 SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_REFERENCE_COUNT_KHR, cl_uint, 2);
629
630 err = clReleaseSemaphoreKHR(sema);
631 test_error(err, "Could not release semaphore");
632 SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_REFERENCE_COUNT_KHR, cl_uint, 1);
633
634 // Confirm that querying CL_SEMAPHORE_PROPERTIES_KHR returns the same
635 // properties the semaphore was created with
636 SEMAPHORE_PARAM_TEST_ARRAY(CL_SEMAPHORE_PROPERTIES_KHR,
637 cl_semaphore_properties_khr, 3, sema_props);
638
639 // Confirm that querying CL_SEMAPHORE_PAYLOAD_KHR returns the unsignaled
640 // state
641 SEMAPHORE_PARAM_TEST(CL_SEMAPHORE_PAYLOAD_KHR, cl_semaphore_payload_khr, 0);
642
643 err = clReleaseSemaphoreKHR(sema);
644 test_error(err, "Could not release semaphore");
645
646 return TEST_PASS;
647 }
648
649 // Test it is possible to export a semaphore to a sync fd and import the same
650 // sync fd to a new semaphore
test_semaphores_import_export_fd(cl_device_id deviceID,cl_context context,cl_command_queue defaultQueue,int num_elements)651 int test_semaphores_import_export_fd(cl_device_id deviceID, cl_context context,
652 cl_command_queue defaultQueue,
653 int num_elements)
654 {
655 cl_int err;
656
657 if (!is_extension_available(deviceID, "cl_khr_semaphore"))
658 {
659 log_info("cl_khr_semaphore is not supported on this platoform. "
660 "Skipping test.\n");
661 return TEST_SKIPPED_ITSELF;
662 }
663
664 if (!is_extension_available(deviceID, "cl_khr_external_semaphore_sync_fd"))
665 {
666 log_info("cl_khr_external_semaphore_sync_fd is not supported on this "
667 "platoform. Skipping test.\n");
668 return TEST_SKIPPED_ITSELF;
669 }
670
671 // Obtain pointers to semaphore's API
672 GET_PFN(deviceID, clCreateSemaphoreWithPropertiesKHR);
673 GET_PFN(deviceID, clEnqueueSignalSemaphoresKHR);
674 GET_PFN(deviceID, clEnqueueWaitSemaphoresKHR);
675 GET_PFN(deviceID, clGetSemaphoreHandleForTypeKHR);
676 GET_PFN(deviceID, clReleaseSemaphoreKHR);
677
678 // Create ooo queue
679 clCommandQueueWrapper queue = clCreateCommandQueue(
680 context, deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
681 test_error(err, "Could not create command queue");
682
683 // Create semaphore
684 cl_semaphore_properties_khr sema_1_props[] = {
685 static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_KHR),
686 static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_BINARY_KHR),
687 static_cast<cl_semaphore_properties_khr>(
688 CL_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR),
689 static_cast<cl_semaphore_properties_khr>(
690 CL_SEMAPHORE_HANDLE_SYNC_FD_KHR),
691 static_cast<cl_semaphore_properties_khr>(
692 CL_SEMAPHORE_EXPORT_HANDLE_TYPES_LIST_END_KHR),
693 0
694 };
695 cl_semaphore_khr sema_1 =
696 clCreateSemaphoreWithPropertiesKHR(context, sema_1_props, &err);
697 test_error(err, "Could not create semaphore");
698
699 // Signal semaphore
700 clEventWrapper signal_event;
701 err = clEnqueueSignalSemaphoresKHR(queue, 1, &sema_1, nullptr, 0, nullptr,
702 &signal_event);
703 test_error(err, "Could not signal semaphore");
704
705 // Extract sync fd
706 int handle = -1;
707 size_t handle_size;
708 err = clGetSemaphoreHandleForTypeKHR(sema_1, deviceID,
709 CL_SEMAPHORE_HANDLE_SYNC_FD_KHR,
710 sizeof(handle), &handle, &handle_size);
711 test_error(err, "Could not extract semaphore handle");
712 test_assert_error(sizeof(handle) == handle_size, "Invalid handle size");
713 test_assert_error(handle >= 0, "Invalid handle");
714
715 // Create semaphore from sync fd
716 cl_semaphore_properties_khr sema_2_props[] = {
717 static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_KHR),
718 static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_BINARY_KHR),
719 CL_SEMAPHORE_HANDLE_SYNC_FD_KHR,
720 static_cast<cl_semaphore_properties_khr>(handle), 0
721 };
722
723 cl_semaphore_khr sema_2 =
724 clCreateSemaphoreWithPropertiesKHR(context, sema_2_props, &err);
725 test_error(err, "Could not create semaphore");
726
727 // Wait semaphore
728 clEventWrapper wait_event;
729 err = clEnqueueWaitSemaphoresKHR(queue, 1, &sema_2, nullptr, 0, nullptr,
730 &wait_event);
731 test_error(err, "Could not wait semaphore");
732
733 // Finish
734 err = clFinish(queue);
735 test_error(err, "Could not finish queue");
736
737 // Check all events are completed
738 test_assert_event_complete(signal_event);
739 test_assert_event_complete(wait_event);
740
741 // Release semaphore
742 err = clReleaseSemaphoreKHR(sema_1);
743 test_error(err, "Could not release semaphore");
744
745 err = clReleaseSemaphoreKHR(sema_2);
746 test_error(err, "Could not release semaphore");
747 return TEST_PASS;
748 }