| /* { dg-do run } */ |
| |
| #include <omp.h> |
| #include <assert.h> |
| #include <unistd.h> // For 'alarm'. |
| |
| #include "on_device_arch.h" |
| |
| /* Test tasks with detach clause on an offload device. Each device |
| thread spawns off a chain of tasks, that can then be executed by |
| any available thread. */ |
| |
| int main (void) |
| { |
| //TODO See '../libgomp.c/pr99555-1.c'. |
| if (on_device_arch_nvptx ()) |
| alarm (4); /*TODO Until resolved, make sure that we exit quickly, with error status. |
| { dg-xfail-run-if "PR99555" { offload_device_nvptx } } */ |
| |
| int x = 0, y = 0, z = 0; |
| int thread_count; |
| omp_event_handle_t detach_event1, detach_event2; |
| |
| #pragma omp target map (tofrom: x, y, z) map (from: thread_count) |
| #pragma omp parallel private (detach_event1, detach_event2) |
| { |
| #pragma omp single |
| thread_count = omp_get_num_threads (); |
| |
| #pragma omp task detach(detach_event1) untied |
| #pragma omp atomic update |
| x++; |
| |
| #pragma omp task detach(detach_event2) untied |
| { |
| #pragma omp atomic update |
| y++; |
| omp_fulfill_event (detach_event1); |
| } |
| |
| #pragma omp task untied |
| { |
| #pragma omp atomic update |
| z++; |
| omp_fulfill_event (detach_event2); |
| } |
| } |
| |
| assert (x == thread_count); |
| assert (y == thread_count); |
| assert (z == thread_count); |
| } |