/* { dg-do run } */

#include <omp.h>
#include <assert.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)
{
  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);
}