4dd9e1c541
... still awaiting proper resolution, of course. libgomp/ PR target/99555 * testsuite/lib/libgomp.exp (check_effective_target_offload_device_nvptx): New. * testsuite/libgomp.c/pr99555-1.c <nvptx offload device>: Until resolved, make sure that we exit quickly, with error status, XFAILed. * testsuite/libgomp.c-c++-common/task-detach-6.c: Likewise. * testsuite/libgomp.fortran/task-detach-6.f90: Likewise.
53 lines
1.2 KiB
C
53 lines
1.2 KiB
C
/* { 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);
|
|
}
|