libgomp/testsuite/libgomp.c-c++-common/task-detach-6.c

   1 /* { dg-do run } */
   2
   3 #include <omp.h>
   4 #include <assert.h>
   5 #include <unistd.h> // For 'alarm'.
   6
   7 #include "on_device_arch.h"
   8
   9 /* Test tasks with detach clause on an offload device.  Each device
  10    thread spawns off a chain of tasks, that can then be executed by
  11    any available thread.  */
  12
  13 int main (void)
  14 {
  15   //TODO See '../libgomp.c/pr99555-1.c'.
  16   if (on_device_arch_nvptx ())
  17     alarm (4); /*TODO Until resolved, make sure that we exit quickly, with error status.
  18                  { dg-xfail-run-if "PR99555" { offload_device_nvptx } } */
  19
  20   int x = 0, y = 0, z = 0;
  21   int thread_count;
  22   omp_event_handle_t detach_event1, detach_event2;
  23
  24   #pragma omp target map (tofrom: x, y, z) map (from: thread_count)
  25     #pragma omp parallel private (detach_event1, detach_event2)
  26       {
  27         #pragma omp single
  28           thread_count = omp_get_num_threads ();
  29
  30         #pragma omp task detach(detach_event1) untied
  31           #pragma omp atomic update
  32             x++;
  33
  34         #pragma omp task detach(detach_event2) untied
  35         {
  36           #pragma omp atomic update
  37             y++;
  38           omp_fulfill_event (detach_event1);
  39         }
  40
  41         #pragma omp task untied
  42         {
  43           #pragma omp atomic update
  44             z++;
  45           omp_fulfill_event (detach_event2);
  46         }
  47       }
  48
  49   assert (x == thread_count);
  50   assert (y == thread_count);
  51   assert (z == thread_count);
  52 }