From patchwork Fri Feb 19 19:12:42 2021
Content-Type: text/plain; charset="utf-8"
MIME-Version: 1.0
Content-Transfer-Encoding: 7bit
Subject: [WIP] Re: [PATCH] openmp: Fix intermittent hanging of task-detach-6
 libgomp tests [PR98738]
X-Patchwork-Submitter: Kwok Cheung Yeung <kcy@codesourcery.com>
X-Patchwork-Id: 49366
Message-Id: <0aca6daf-356a-7b03-c007-e33c35114356@codesourcery.com>
To: Jakub Jelinek <jakub@redhat.com>
Cc: GCC Patches <gcc-patches@gcc.gnu.org>
Date: Fri, 19 Feb 2021 19:12:42 +0000
From: Kwok Cheung Yeung <kcy@codesourcery.com>
List-Id: Gcc-patches mailing list <gcc-patches.gcc.gnu.org>

Hello

Sorry for taking so long in replying.

On 29/01/2021 3:03 pm, Jakub Jelinek wrote:
> It can also crash if team is NULL, which will happen any time
> this is called outside of a parallel.  Just try (should go into testsuite
> too):
> #include <omp.h>
> 
> int
> main ()
> {
>    omp_event_handle_t ev;
>    #pragma omp task detach (ev)
>    omp_fulfill_event (ev);
>    return 0;
> }
>

I have included this as task-detach-11.{c|f90}.

> Additionally, there is an important difference between fulfill for
> included tasks and for non-included tasks, for the former there is no team
> or anything to care about, for the latter there is a team and one needs to
> take the task_lock, but at that point it can do pretty much everything in
> omp_fulfill_event rather than handling it elsewhere.
> 
> So, what I'm suggesting is:
> 
> Replace
>    bool detach;
>    gomp_sem_t completion_sem;
> with
>    struct gomp_task_detach *detach;
> and add struct gomp_task_detach that would contain everything that will be
> needed (indirect so that we don't waste space for it in every task, but only
> for those that have detach clause).
> We need:
> 1) some way to tell if it is an included task or not
> 2) for included tasks the gomp_sem_t completion_sem
> (and nothing but 1) and 2) for those),
> 3) struct gomp_team * for non-included tasks
> 4) some way to find out if the task has finished and is just waiting for
> fulfill event (perhaps your GOMP_TASK_DETACHED is ok for that)
> 5) some way to find out if the task has been fulfilled already
> (gomp_sem_t for that seems an overkill though)
> 
> 1) could be done through the struct gomp_team *team; member,
> set it to NULL in included tasks (no matter if they are in some team or not)
> and to non-NULL team of the task (non-included tasks must have a team).
> 

I have opted for a union of completion_sem (for tasks that are undeferred) and a 
struct gomp_team *detach_team (for deferred tasks) that holds the team if the 
completion event has not yet fulfilled, or NULL if is it. I don't see the point 
of having an indirection to the union here since the union is just the size of a 
pointer, so it might as well be inlined.

> And I don't see the point of task_detach_queue if we can handle the
> dependers etc. all in omp_fulfill_event, which I think we can if we take the
> task_lock.

I have removed the task_detach_queue. The team barrier, taskwait and 
taskgroup_end now just set the task kind to GOMP_TASK_DETACHED without 
decrementing the task_count if a task finishes with detach_team non-NULL.

> So, I think omp_fulfill_event should look at the task->detach it got,
> if task->detach->team is NULL, it is included task, GOMP_task should have
> initialized task->detach->completion_sem and omp_fulfill_event should just
> gomp_sem_post it and that is all, GOMP_task for included task needs to
> gomp_sem_wait after it finishes before it returns.

omp_fulfill_event now posts completion_sem if the task kind is 
OMP_TASK_UNDEFERRED, and GOMP_task waits for it. Since the task is executed 
within GOMP_task, it already knows if the task has a detach clause or not, so we 
do not need to store that information in gomp_task.

> Otherwise, take the team's task_lock, and look at whether the task is still
> running, in that case just set the bool that it has been fulfilled (or
> whatever way of signalling 5), perhaps it can be say clearing task->detach
> pointer).

detach_team is now set to NULL when the event is fulfilled if the task has not 
started yet or is still executing (checked by the kind). In that case, when the 
task finishes executing, it behaves just like a task without detach would and 
finishes normally.

   When creating non-included tasks in GOMP_task with detach clause
> through gomp_malloc, it would add the size needed for struct
> gomp_task_detach.

Not necessary with the inlined union.

> But if the task is already in GOMP_TASK_DETACHED state, instead we need
> while holding the task_lock do everything that would have been done normally
> on task finish, but we've skipped it because it hasn't been fulfilled.
> Including the waking/sem_posts when something could be waiting on that task.
> 
> Do you agree with this, or see some reason why this can't work?

The main problem I see is this code in gomp_barrier_handle_tasks:

	  if (--team->task_count == 0
	      && gomp_team_barrier_waiting_for_tasks (&team->barrier))
	    {
	      gomp_team_barrier_done (&team->barrier, state);

We do not have access to state from within omp_fulfill_event, so how should this 
be handled?

> And testsuite should include also cases where we wait for the tasks with
> detach clause to be fulfilled at the end of taskgroup (i.e. need to cover
> all of taskwait, taskgroup end and barrier).

I have changed task-detach-[56].* to test the barrier, task-detach-[78].* to 
test taskwait, and task-detach-(9|10) to test taskgroup (with the first one 
without a target construct, the second with).

I have included the current state of my patch. All task-detach-* tests pass when 
executed without offloading or with offloading to GCN, but with offloading to 
Nvidia, task-detach-6.* hangs consistently but everything else passes (probably 
because of the missing gomp_team_barrier_done?).

Kwok
>From 31a5c736910036364fd1f0f3cf7ac28437864a27 Mon Sep 17 00:00:00 2001
From: Kwok Cheung Yeung <kcy@codesourcery.com>
Date: Thu, 21 Jan 2021 05:38:47 -0800
Subject: [PATCH] openmp: Fix intermittent hanging of task-detach-6 libgomp
 tests [PR98738]

This adds support for the task detach clause to taskwait and taskgroup, and
simplifies the handling of the detach clause by moving most of the extra
handling required for detach tasks to omp_fulfill_event.

2021-02-19  Kwok Cheung Yeung  <kcy@codesourcery.com>

	libgomp/

	PR libgomp/98738
	* libgomp.h (enum gomp_task_kind): Add GOMP_TASK_DETACHED.
	(struct gomp_task): Replace detach and completion_sem fields with
	union containing completion_sem and detach_team.
	(struct gomp_team): Remove task_detach_queue.
	* task.c: Include assert.h.
	(gomp_init_task): Initialize detach_team field.
	(task_fulfilled_p): Delete.
	(GOMP_task): Use address of task as the event handle.  Remove
	initialization of detach field.  Initialize detach_team field for
	deferred tasks.
	(gomp_barrier_handle_tasks): Remove handling of task_detach_queue.
	Set kind of suspended detach task to GOMP_TASK_DETACHED and
	decrement task_running_count.  Move finish_cancelled block out of
	else branch.
	(GOMP_taskwait): Handle tasks with completion events that have not
	been fulfilled.
	(GOMP_taskgroup_end): Likewise.
	(omp_fulfill_event): Use address of task as event handle.  Post to
	completion_sem for undeferred tasks.  Clear detach_team if task
	has not finished.  For finished tasks, handle post-execution tasks,
	post to taskwait_sem and taskgroup_sem if necessary, call
	gomp_team_barrier_wake if necessary, and free task.
	* testsuite/libgomp.c-c++-common/task-detach-1.c: Fix formatting.
	* testsuite/libgomp.c-c++-common/task-detach-2.c: Fix formatting.
	* testsuite/libgomp.c-c++-common/task-detach-3.c: Fix formatting.
	* testsuite/libgomp.c-c++-common/task-detach-4.c: Fix formatting.
	* testsuite/libgomp.c-c++-common/task-detach-5.c: Fix formatting.
	Change data-sharing of detach events on enclosing parallel to private.
	* testsuite/libgomp.c-c++-common/task-detach-6.c: Likewise.  Remove
	taskwait directive.
	* testsuite/libgomp.c-c++-common/task-detach-7.c: New.
	* testsuite/libgomp.c-c++-common/task-detach-8.c: New.
	* testsuite/libgomp.c-c++-common/task-detach-9.c: New.
	* testsuite/libgomp.c-c++-common/task-detach-10.c: New.
	* testsuite/libgomp.c-c++-common/task-detach-11.c: New.
	* testsuite/libgomp.c-c++-common/task-detach-1.f90: Fix formatting.
	* testsuite/libgomp.c-c++-common/task-detach-2.f90: Fix formatting.
	* testsuite/libgomp.c-c++-common/task-detach-3.f90: Fix formatting.
	* testsuite/libgomp.c-c++-common/task-detach-4.f90: Fix formatting.
	* testsuite/libgomp.fortran/task-detach-5.f90: Fix formatting.
	Change data-sharing of detach events on enclosing parallel to private.
	* testsuite/libgomp.fortran/task-detach-6.f90: Likewise.  Remove
	taskwait directive.
	* testsuite/libgomp.c-c++-common/task-detach-7.f90: New.
	* testsuite/libgomp.c-c++-common/task-detach-8.f90: New.
	* testsuite/libgomp.c-c++-common/task-detach-9.f90: New.
	* testsuite/libgomp.c-c++-common/task-detach-10.f90: New.
	* testsuite/libgomp.c-c++-common/task-detach-11.f90: New.
---
 libgomp/libgomp.h                                  |  18 +-
 libgomp/task.c                                     | 225 +++++++++++++--------
 libgomp/team.c                                     |   2 -
 .../testsuite/libgomp.c-c++-common/task-detach-1.c |   4 +-
 .../libgomp.c-c++-common/task-detach-10.c          |  45 +++++
 .../libgomp.c-c++-common/task-detach-11.c          |  13 ++
 .../testsuite/libgomp.c-c++-common/task-detach-2.c |   6 +-
 .../testsuite/libgomp.c-c++-common/task-detach-3.c |   6 +-
 .../testsuite/libgomp.c-c++-common/task-detach-4.c |   4 +-
 .../testsuite/libgomp.c-c++-common/task-detach-5.c |   8 +-
 .../testsuite/libgomp.c-c++-common/task-detach-6.c |   8 +-
 .../testsuite/libgomp.c-c++-common/task-detach-7.c |  45 +++++
 .../testsuite/libgomp.c-c++-common/task-detach-8.c |  47 +++++
 .../testsuite/libgomp.c-c++-common/task-detach-9.c |  43 ++++
 .../testsuite/libgomp.fortran/task-detach-1.f90    |   4 +-
 .../testsuite/libgomp.fortran/task-detach-10.f90   |  44 ++++
 .../testsuite/libgomp.fortran/task-detach-11.f90   |  13 ++
 .../testsuite/libgomp.fortran/task-detach-2.f90    |   6 +-
 .../testsuite/libgomp.fortran/task-detach-3.f90    |   6 +-
 .../testsuite/libgomp.fortran/task-detach-4.f90    |   4 +-
 .../testsuite/libgomp.fortran/task-detach-5.f90    |   8 +-
 .../testsuite/libgomp.fortran/task-detach-6.f90    |  16 +-
 .../testsuite/libgomp.fortran/task-detach-7.f90    |  42 ++++
 .../testsuite/libgomp.fortran/task-detach-8.f90    |  45 +++++
 .../testsuite/libgomp.fortran/task-detach-9.f90    |  41 ++++
 25 files changed, 573 insertions(+), 130 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/task-detach-10.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/task-detach-11.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/task-detach-7.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/task-detach-8.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/task-detach-9.c
 create mode 100644 libgomp/testsuite/libgomp.fortran/task-detach-10.f90
 create mode 100644 libgomp/testsuite/libgomp.fortran/task-detach-11.f90
 create mode 100644 libgomp/testsuite/libgomp.fortran/task-detach-7.f90
 create mode 100644 libgomp/testsuite/libgomp.fortran/task-detach-8.f90
 create mode 100644 libgomp/testsuite/libgomp.fortran/task-detach-9.f90

diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index b4d0c93..90a6f02 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -481,7 +481,10 @@ enum gomp_task_kind
      but not yet completed.  Once that completes, they will be readded
      into the queues as GOMP_TASK_WAITING in order to perform the var
      unmapping.  */
-  GOMP_TASK_ASYNC_RUNNING
+  GOMP_TASK_ASYNC_RUNNING,
+  /* Task that has finished executing but is waiting for its
+     completion event to be fulfilled.  */
+  GOMP_TASK_DETACHED
 };
 
 struct gomp_task_depend_entry
@@ -545,8 +548,14 @@ struct gomp_task
      entries and the gomp_task in which they reside.  */
   struct priority_node pnode[3];
 
-  bool detach;
-  gomp_sem_t completion_sem;
+  union {
+    /* Valid only if kind == GOMP_TASK_UNDEFERRED.  */
+    gomp_sem_t completion_sem;
+    /* Valid for other values of kind.  Set to the team that executes the
+       task if the task is detached and the completion event has yet to be
+       fulfilled.  */
+    struct gomp_team *detach_team;
+  };
 
   struct gomp_task_icv icv;
   void (*fn) (void *);
@@ -688,8 +697,7 @@ struct gomp_team
   int work_share_cancelled;
   int team_cancelled;
 
-  /* Tasks waiting for their completion event to be fulfilled.  */
-  struct priority_queue task_detach_queue;
+  /* Number of tasks waiting for their completion event to be fulfilled.  */
   unsigned int task_detach_count;
 
   /* This array contains structures for implicit tasks.  */
diff --git a/libgomp/task.c b/libgomp/task.c
index b242e7c..399e18b 100644
--- a/libgomp/task.c
+++ b/libgomp/task.c
@@ -29,6 +29,7 @@
 #include "libgomp.h"
 #include <stdlib.h>
 #include <string.h>
+#include <assert.h>
 #include "gomp-constants.h"
 
 typedef struct gomp_task_depend_entry *hash_entry_type;
@@ -86,7 +87,7 @@ gomp_init_task (struct gomp_task *task, struct gomp_task *parent_task,
   task->dependers = NULL;
   task->depend_hash = NULL;
   task->depend_count = 0;
-  task->detach = false;
+  task->detach_team = NULL;
 }
 
 /* Clean up a task, after completing it.  */
@@ -327,12 +328,6 @@ gomp_task_handle_depend (struct gomp_task *task, struct gomp_task *parent,
     }
 }
 
-static bool
-task_fulfilled_p (struct gomp_task *task)
-{
-  return gomp_sem_getcount (&task->completion_sem) > 0;
-}
-
 /* Called when encountering an explicit task directive.  If IF_CLAUSE is
    false, then we must not delay in executing the task.  If UNTIED is true,
    then the task may be executed by any member of the team.
@@ -417,13 +412,13 @@ GOMP_task (void (*fn) (void *), void *data, void (*cpyfn) (void *, void *),
 
       if ((flags & GOMP_TASK_FLAG_DETACH) != 0)
 	{
-	  task.detach = true;
 	  gomp_sem_init (&task.completion_sem, 0);
-	  *(void **) detach = &task.completion_sem;
+	  *(void **) detach = &task;
 	  if (data)
-	    *(void **) data = &task.completion_sem;
+	    *(void **) data = &task;
 
-	  gomp_debug (0, "New event: %p\n", &task.completion_sem);
+	  gomp_debug (0, "Thread %d: new event: %p\n",
+		      thr->ts.team_id, &task);
 	}
 
       if (thr->task)
@@ -443,7 +438,7 @@ GOMP_task (void (*fn) (void *), void *data, void (*cpyfn) (void *, void *),
       else
 	fn (data);
 
-      if (task.detach && !task_fulfilled_p (&task))
+      if ((flags & GOMP_TASK_FLAG_DETACH) != 0 && detach)
 	gomp_sem_wait (&task.completion_sem);
 
       /* Access to "children" is normally done inside a task_lock
@@ -481,18 +476,17 @@ GOMP_task (void (*fn) (void *), void *data, void (*cpyfn) (void *, void *),
 		      & ~(uintptr_t) (arg_align - 1));
       gomp_init_task (task, parent, gomp_icv (false));
       task->priority = priority;
-      task->kind = GOMP_TASK_UNDEFERRED;
       task->in_tied_task = parent->in_tied_task;
       task->taskgroup = taskgroup;
       if ((flags & GOMP_TASK_FLAG_DETACH) != 0)
 	{
-	  task->detach = true;
-	  gomp_sem_init (&task->completion_sem, 0);
-	  *(void **) detach = &task->completion_sem;
+	  task->detach_team = team;
+
+	  *(void **) detach = task;
 	  if (data)
-	    *(void **) data = &task->completion_sem;
+	    *(void **) data = task;
 
-	  gomp_debug (0, "New event: %p\n", &task->completion_sem);
+	  gomp_debug (0, "Thread %d: new event: %p\n", thr->ts.team_id, task);
 	}
       thr->task = task;
       if (cpyfn)
@@ -1362,27 +1356,6 @@ gomp_barrier_handle_tasks (gomp_barrier_state_t state)
     {
       bool cancelled = false;
 
-      /* Look for a queued detached task with a fulfilled completion event
-	 that is ready to finish.  */
-      child_task = priority_queue_find (PQ_TEAM, &team->task_detach_queue,
-					task_fulfilled_p);
-      if (child_task)
-	{
-	  priority_queue_remove (PQ_TEAM, &team->task_detach_queue,
-				 child_task, MEMMODEL_RELAXED);
-	  --team->task_detach_count;
-	  gomp_debug (0, "thread %d: found task with fulfilled event %p\n",
-		      thr->ts.team_id, &child_task->completion_sem);
-
-	if (to_free)
-	  {
-	    gomp_finish_task (to_free);
-	    free (to_free);
-	    to_free = NULL;
-	  }
-	  goto finish_cancelled;
-	}
-
       if (!priority_queue_empty_p (&team->task_queue, MEMMODEL_RELAXED))
 	{
 	  bool ignored;
@@ -1450,43 +1423,43 @@ gomp_barrier_handle_tasks (gomp_barrier_state_t state)
       gomp_mutex_lock (&team->task_lock);
       if (child_task)
 	{
-	  if (child_task->detach && !task_fulfilled_p (child_task))
+	  if (child_task->detach_team)
 	    {
-	      priority_queue_insert (PQ_TEAM, &team->task_detach_queue,
-				     child_task, child_task->priority,
-				     PRIORITY_INSERT_END,
-				     false, false);
+	      assert (child_task->detach_team == team);
+	      child_task->kind = GOMP_TASK_DETACHED;
 	      ++team->task_detach_count;
-	      gomp_debug (0, "thread %d: queueing task with event %p\n",
-			  thr->ts.team_id, &child_task->completion_sem);
+	      --team->task_running_count;
+	      gomp_debug (0,
+			  "thread %d: task with event %p finished without "
+			  "completion event fulfilled in team barrier\n",
+			  thr->ts.team_id, child_task);
 	      child_task = NULL;
+	      continue;
 	    }
-	  else
+
+	 finish_cancelled:;
+	  size_t new_tasks
+	    = gomp_task_run_post_handle_depend (child_task, team);
+	  gomp_task_run_post_remove_parent (child_task);
+	  gomp_clear_parent (&child_task->children_queue);
+	  gomp_task_run_post_remove_taskgroup (child_task);
+	  to_free = child_task;
+	  if (!cancelled)
+	    team->task_running_count--;
+	  child_task = NULL;
+	  if (new_tasks > 1)
 	    {
-	     finish_cancelled:;
-	      size_t new_tasks
-		= gomp_task_run_post_handle_depend (child_task, team);
-	      gomp_task_run_post_remove_parent (child_task);
-	      gomp_clear_parent (&child_task->children_queue);
-	      gomp_task_run_post_remove_taskgroup (child_task);
-	      to_free = child_task;
-	      child_task = NULL;
-	      if (!cancelled)
-		team->task_running_count--;
-	      if (new_tasks > 1)
-		{
-		  do_wake = team->nthreads - team->task_running_count;
-		  if (do_wake > new_tasks)
-		    do_wake = new_tasks;
-		}
-	      if (--team->task_count == 0
-		  && gomp_team_barrier_waiting_for_tasks (&team->barrier))
-		{
-		  gomp_team_barrier_done (&team->barrier, state);
-		  gomp_mutex_unlock (&team->task_lock);
-		  gomp_team_barrier_wake (&team->barrier, 0);
-		  gomp_mutex_lock (&team->task_lock);
-		}
+	      do_wake = team->nthreads - team->task_running_count;
+	      if (do_wake > new_tasks)
+		do_wake = new_tasks;
+	    }
+	  if (--team->task_count == 0
+	      && gomp_team_barrier_waiting_for_tasks (&team->barrier))
+	    {
+	      gomp_team_barrier_done (&team->barrier, state);
+	      gomp_mutex_unlock (&team->task_lock);
+	      gomp_team_barrier_wake (&team->barrier, 0);
+	      gomp_mutex_lock (&team->task_lock);
 	    }
 	}
     }
@@ -1559,7 +1532,8 @@ GOMP_taskwait (void)
       else
 	{
 	/* All tasks we are waiting for are either running in other
-	   threads, or they are tasks that have not had their
+	   threads, are detached and waiting for the completion event to be
+	   fulfilled, or they are tasks that have not had their
 	   dependencies met (so they're not even in the queue).  Wait
 	   for them.  */
 	  if (task->taskwait == NULL)
@@ -1614,6 +1588,19 @@ GOMP_taskwait (void)
       gomp_mutex_lock (&team->task_lock);
       if (child_task)
 	{
+	  if (child_task->detach_team)
+	    {
+	      assert (child_task->detach_team == team);
+	      child_task->kind = GOMP_TASK_DETACHED;
+	      ++team->task_detach_count;
+	      gomp_debug (0,
+			  "thread %d: task with event %p finished without "
+			  "completion event fulfilled in taskwait\n",
+			  thr->ts.team_id, child_task);
+	      child_task = NULL;
+	      continue;
+	    }
+
 	 finish_cancelled:;
 	  size_t new_tasks
 	    = gomp_task_run_post_handle_depend (child_task, team);
@@ -2069,6 +2056,19 @@ GOMP_taskgroup_end (void)
       gomp_mutex_lock (&team->task_lock);
       if (child_task)
 	{
+	  if (child_task->detach_team)
+	    {
+	      assert (child_task->detach_team == team);
+	      child_task->kind = GOMP_TASK_DETACHED;
+	      ++team->task_detach_count;
+	      gomp_debug (0,
+			  "thread %d: task with event %p finished without "
+			  "completion event fulfilled in taskgroup\n",
+			  thr->ts.team_id, child_task);
+	      child_task = NULL;
+	      continue;
+	    }
+
 	 finish_cancelled:;
 	  size_t new_tasks
 	    = gomp_task_run_post_handle_depend (child_task, team);
@@ -2402,17 +2402,80 @@ ialias (omp_in_final)
 void
 omp_fulfill_event (omp_event_handle_t event)
 {
-  gomp_sem_t *sem = (gomp_sem_t *) event;
-  struct gomp_thread *thr = gomp_thread ();
-  struct gomp_team *team = thr ? thr->ts.team : NULL;
+  struct gomp_task *task = (struct gomp_task *) event;
+  if (task->kind == GOMP_TASK_UNDEFERRED)
+  {
+    if (gomp_sem_getcount (&task->completion_sem) > 0)
+      gomp_fatal ("omp_fulfill_event: %p event already fulfilled!\n", task);
+
+    gomp_debug (0, "omp_fulfill_event: %p event for undeferred task\n", task);
+    gomp_sem_post (&task->completion_sem);
+    return;
+  }
 
-  if (gomp_sem_getcount (sem) > 0)
-    gomp_fatal ("omp_fulfill_event: %p event already fulfilled!\n", sem);
+  struct gomp_team *team = task->detach_team;
+  if (!team)
+    gomp_fatal ("omp_fulfill_event: %p event is invalid or has already "
+		"been fulfilled!\n", task);
 
-  gomp_debug (0, "omp_fulfill_event: %p\n", sem);
-  gomp_sem_post (sem);
-  if (team)
-    gomp_team_barrier_wake (&team->barrier, 1);
+  gomp_mutex_lock (&team->task_lock);
+  if (task->kind != GOMP_TASK_DETACHED)
+    {
+      /* The task has not finished running yet.  */
+      gomp_debug (0,
+		  "omp_fulfill_event: %p event fulfilled for unfinished "
+		  "task\n", task);
+      task->detach_team = NULL;
+      gomp_mutex_unlock (&team->task_lock);
+      return;
+    }
+
+  gomp_debug (0, "omp_fulfill_event: %p event fulfilled for finished task\n",
+	      task);
+  size_t new_tasks = gomp_task_run_post_handle_depend (task, team);
+  gomp_task_run_post_remove_parent (task);
+  gomp_clear_parent (&task->children_queue);
+  gomp_task_run_post_remove_taskgroup (task);
+  team->task_count--;
+  team->task_detach_count--;
+
+  /* Wake up any threads that may be waiting for the detached task
+     to complete.  */
+  struct gomp_task *parent = task->parent;
+
+  if (parent && parent->taskwait)
+    {
+      if (parent->taskwait->in_taskwait)
+	{
+	  parent->taskwait->in_taskwait = false;
+	  gomp_sem_post (&parent->taskwait->taskwait_sem);
+	}
+      else if (parent->taskwait->in_depend_wait)
+	{
+	  parent->taskwait->in_depend_wait = false;
+	  gomp_sem_post (&parent->taskwait->taskwait_sem);
+	}
+    }
+  if (task->taskgroup && task->taskgroup->in_taskgroup_wait)
+    {
+      task->taskgroup->in_taskgroup_wait = false;
+      gomp_sem_post (&task->taskgroup->taskgroup_sem);
+    }
+
+  int do_wake = 0;
+  if (new_tasks > 1)
+    {
+      do_wake = team->nthreads - team->task_running_count;
+      if (do_wake > new_tasks)
+	do_wake = new_tasks;
+    }
+
+  gomp_mutex_unlock (&team->task_lock);
+  if (do_wake)
+    gomp_team_barrier_wake (&team->barrier, do_wake);
+
+  gomp_finish_task (task);
+  free (task);
 }
 
 ialias (omp_fulfill_event)
diff --git a/libgomp/team.c b/libgomp/team.c
index 0f3707c..9662234 100644
--- a/libgomp/team.c
+++ b/libgomp/team.c
@@ -206,7 +206,6 @@ gomp_new_team (unsigned nthreads)
   team->work_share_cancelled = 0;
   team->team_cancelled = 0;
 
-  priority_queue_init (&team->task_detach_queue);
   team->task_detach_count = 0;
 
   return team;
@@ -224,7 +223,6 @@ free_team (struct gomp_team *team)
   gomp_barrier_destroy (&team->barrier);
   gomp_mutex_destroy (&team->task_lock);
   priority_queue_free (&team->task_queue);
-  priority_queue_free (&team->task_detach_queue);
   team_free (team);
 }
 
diff --git a/libgomp/testsuite/libgomp.c-c++-common/task-detach-1.c b/libgomp/testsuite/libgomp.c-c++-common/task-detach-1.c
index 8583e37..14932b0 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/task-detach-1.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-1.c
@@ -14,10 +14,10 @@ int main (void)
   #pragma omp parallel
     #pragma omp single
     {
-      #pragma omp task detach(detach_event1)
+      #pragma omp task detach (detach_event1)
 	x++;
 
-      #pragma omp task detach(detach_event2)
+      #pragma omp task detach (detach_event2)
       {
 	y++;
 	omp_fulfill_event (detach_event1);
diff --git a/libgomp/testsuite/libgomp.c-c++-common/task-detach-10.c b/libgomp/testsuite/libgomp.c-c++-common/task-detach-10.c
new file mode 100644
index 0000000..10d6746
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-10.c
@@ -0,0 +1,45 @@
+/* { 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 in a taskgroup, 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 taskgroup
+	{
+	  #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);
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/task-detach-11.c b/libgomp/testsuite/libgomp.c-c++-common/task-detach-11.c
new file mode 100644
index 0000000..dd002dc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-11.c
@@ -0,0 +1,13 @@
+/* { dg-do run } */
+
+#include <omp.h>
+
+/* Test the detach clause when the task is undeferred.  */
+
+int main (void)
+{
+  omp_event_handle_t event;
+
+  #pragma omp task detach (event)
+    omp_fulfill_event (event);
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/task-detach-2.c b/libgomp/testsuite/libgomp.c-c++-common/task-detach-2.c
index 943ac2a..3e33c40 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/task-detach-2.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-2.c
@@ -12,13 +12,13 @@ int main (void)
   omp_event_handle_t detach_event1, detach_event2;
   int x = 0, y = 0, z = 0;
 
-  #pragma omp parallel num_threads(1)
+  #pragma omp parallel num_threads (1)
     #pragma omp single
     {
-      #pragma omp task detach(detach_event1)
+      #pragma omp task detach (detach_event1)
 	x++;
 
-      #pragma omp task detach(detach_event2)
+      #pragma omp task detach (detach_event2)
       {
 	y++;
 	omp_fulfill_event (detach_event1);
diff --git a/libgomp/testsuite/libgomp.c-c++-common/task-detach-3.c b/libgomp/testsuite/libgomp.c-c++-common/task-detach-3.c
index 2609fb1..c85857d 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/task-detach-3.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-3.c
@@ -14,16 +14,16 @@ int main (void)
   #pragma omp parallel
     #pragma omp single
     {
-      #pragma omp task depend(out:dep) detach(detach_event)
+      #pragma omp task depend (out:dep) detach (detach_event)
 	x++;
 
       #pragma omp task
       {
 	y++;
-	omp_fulfill_event(detach_event);
+	omp_fulfill_event (detach_event);
       }
 
-      #pragma omp task depend(in:dep)
+      #pragma omp task depend (in:dep)
 	z++;
     }
 
diff --git a/libgomp/testsuite/libgomp.c-c++-common/task-detach-4.c b/libgomp/testsuite/libgomp.c-c++-common/task-detach-4.c
index eeb9554..cd0d2b3 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/task-detach-4.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-4.c
@@ -14,10 +14,10 @@ int main (void)
 
   #pragma omp parallel
     #pragma omp single
-      #pragma omp task detach(detach_event)
+      #pragma omp task detach (detach_event)
       {
 	x++;
-	omp_fulfill_event(detach_event);
+	omp_fulfill_event (detach_event);
       }
 
   assert (x == 1);
diff --git a/libgomp/testsuite/libgomp.c-c++-common/task-detach-5.c b/libgomp/testsuite/libgomp.c-c++-common/task-detach-5.c
index 5a01517..382f377 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/task-detach-5.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-5.c
@@ -12,16 +12,16 @@ int main (void)
   int thread_count;
   omp_event_handle_t detach_event1, detach_event2;
 
-  #pragma omp parallel firstprivate(detach_event1, detach_event2)
+  #pragma omp parallel private (detach_event1, detach_event2)
   {
     #pragma omp single
-      thread_count = omp_get_num_threads();
+      thread_count = omp_get_num_threads ();
 
-    #pragma omp task detach(detach_event1) untied
+    #pragma omp task detach (detach_event1) untied
       #pragma omp atomic update
 	x++;
 
-    #pragma omp task detach(detach_event2) untied
+    #pragma omp task detach (detach_event2) untied
     {
       #pragma omp atomic update
 	y++;
diff --git a/libgomp/testsuite/libgomp.c-c++-common/task-detach-6.c b/libgomp/testsuite/libgomp.c-c++-common/task-detach-6.c
index b5f68cc..e5c2291 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/task-detach-6.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-6.c
@@ -13,11 +13,11 @@ int main (void)
   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 firstprivate(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();
+	  thread_count = omp_get_num_threads ();
 
 	#pragma omp task detach(detach_event1) untied
 	  #pragma omp atomic update
@@ -36,8 +36,6 @@ int main (void)
 	    z++;
 	  omp_fulfill_event (detach_event2);
 	}
-
-	#pragma omp taskwait
       }
 
   assert (x == thread_count);
diff --git a/libgomp/testsuite/libgomp.c-c++-common/task-detach-7.c b/libgomp/testsuite/libgomp.c-c++-common/task-detach-7.c
new file mode 100644
index 0000000..3f025d6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-7.c
@@ -0,0 +1,45 @@
+/* { dg-do run } */
+
+#include <omp.h>
+#include <assert.h>
+
+/* Test tasks with detach clause.  Each thread spawns off a chain of tasks,
+   that can then be executed by any available thread.  Each thread uses
+   taskwait to wait for the child tasks to complete.  */
+
+int main (void)
+{
+  int x = 0, y = 0, z = 0;
+  int thread_count;
+  omp_event_handle_t detach_event1, detach_event2;
+
+  #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);
+    }
+
+    #pragma omp taskwait
+  }
+
+  assert (x == thread_count);
+  assert (y == thread_count);
+  assert (z == thread_count);
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/task-detach-8.c b/libgomp/testsuite/libgomp.c-c++-common/task-detach-8.c
new file mode 100644
index 0000000..6f77f12
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-8.c
@@ -0,0 +1,47 @@
+/* { 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.  Each thread uses taskwait to wait for the
+   child tasks to complete.  */
+
+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);
+	}
+
+	#pragma omp taskwait
+      }
+
+  assert (x == thread_count);
+  assert (y == thread_count);
+  assert (z == thread_count);
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/task-detach-9.c b/libgomp/testsuite/libgomp.c-c++-common/task-detach-9.c
new file mode 100644
index 0000000..5316ca5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-9.c
@@ -0,0 +1,43 @@
+/* { dg-do run } */
+
+#include <omp.h>
+#include <assert.h>
+
+/* Test tasks with detach clause.  Each thread spawns off a chain of tasks
+   in a taskgroup, 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 parallel private (detach_event1, detach_event2)
+    #pragma omp taskgroup
+    {
+      #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);
+}
diff --git a/libgomp/testsuite/libgomp.fortran/task-detach-1.f90 b/libgomp/testsuite/libgomp.fortran/task-detach-1.f90
index 217bf65..c53b1ca 100644
--- a/libgomp/testsuite/libgomp.fortran/task-detach-1.f90
+++ b/libgomp/testsuite/libgomp.fortran/task-detach-1.f90
@@ -11,11 +11,11 @@ program task_detach_1
 
   !$omp parallel
     !$omp single
-      !$omp task detach(detach_event1)
+      !$omp task detach (detach_event1)
         x = x + 1
       !$omp end task
 
-      !$omp task detach(detach_event2)
+      !$omp task detach (detach_event2)
         y = y + 1
 	call omp_fulfill_event (detach_event1)
       !$omp end task
diff --git a/libgomp/testsuite/libgomp.fortran/task-detach-10.f90 b/libgomp/testsuite/libgomp.fortran/task-detach-10.f90
new file mode 100644
index 0000000..61f0ea8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/task-detach-10.f90
@@ -0,0 +1,44 @@
+! { dg-do run }
+
+! Test tasks with detach clause on an offload device.  Each device
+! thread spawns off a chain of tasks in a taskgroup, that can then
+! be executed by any available thread.
+
+program task_detach_10
+  use omp_lib
+
+  integer (kind=omp_event_handle_kind) :: detach_event1, detach_event2
+  integer :: x = 0, y = 0, z = 0
+  integer :: thread_count
+
+  !$omp target map (tofrom: x, y, z) map (from: thread_count)
+    !$omp parallel private (detach_event1, detach_event2)
+      !$omp taskgroup
+	!$omp single
+	  thread_count = omp_get_num_threads ()
+	!$omp end single
+
+	!$omp task detach (detach_event1) untied
+	  !$omp atomic update
+	    x = x + 1
+	!$omp end task
+
+	!$omp task detach (detach_event2) untied
+	  !$omp atomic update
+	    y = y + 1
+	  call omp_fulfill_event (detach_event1)
+	!$omp end task
+
+	!$omp task untied
+	  !$omp atomic update
+	    z = z + 1
+	  call omp_fulfill_event (detach_event2)
+	!$omp end task
+      !$omp end taskgroup
+    !$omp end parallel
+  !$omp end target
+
+  if (x /= thread_count) stop 1
+  if (y /= thread_count) stop 2
+  if (z /= thread_count) stop 3
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/task-detach-11.f90 b/libgomp/testsuite/libgomp.fortran/task-detach-11.f90
new file mode 100644
index 0000000..b33baff
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/task-detach-11.f90
@@ -0,0 +1,13 @@
+! { dg-do run }
+
+! Test the detach clause when the task is undeferred.
+
+program task_detach_11
+  use omp_lib
+
+  integer (kind=omp_event_handle_kind) :: detach_event
+
+  !$omp task detach (detach_event)
+    call omp_fulfill_event (detach_event)
+  !$omp end task
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/task-detach-2.f90 b/libgomp/testsuite/libgomp.fortran/task-detach-2.f90
index ecb4829..68e3ff2 100644
--- a/libgomp/testsuite/libgomp.fortran/task-detach-2.f90
+++ b/libgomp/testsuite/libgomp.fortran/task-detach-2.f90
@@ -10,13 +10,13 @@ program task_detach_2
   integer (kind=omp_event_handle_kind) :: detach_event1, detach_event2
   integer :: x = 0, y = 0, z = 0
 
-  !$omp parallel num_threads(1)
+  !$omp parallel num_threads (1)
     !$omp single
-      !$omp task detach(detach_event1)
+      !$omp task detach (detach_event1)
         x = x + 1
       !$omp end task
 
-      !$omp task detach(detach_event2)
+      !$omp task detach (detach_event2)
         y = y + 1
 	call omp_fulfill_event (detach_event1)
       !$omp end task
diff --git a/libgomp/testsuite/libgomp.fortran/task-detach-3.f90 b/libgomp/testsuite/libgomp.fortran/task-detach-3.f90
index bdf93a5..5ac68d5 100644
--- a/libgomp/testsuite/libgomp.fortran/task-detach-3.f90
+++ b/libgomp/testsuite/libgomp.fortran/task-detach-3.f90
@@ -12,16 +12,16 @@ program task_detach_3
 
   !$omp parallel
     !$omp single
-      !$omp task depend(out:dep) detach(detach_event)
+      !$omp task depend (out:dep) detach (detach_event)
         x = x + 1
       !$omp end task
 
       !$omp task
         y = y + 1
-	call omp_fulfill_event(detach_event)
+	call omp_fulfill_event (detach_event)
       !$omp end task
 
-      !$omp task depend(in:dep)
+      !$omp task depend (in:dep)
         z = z + 1
       !$omp end task
     !$omp end single
diff --git a/libgomp/testsuite/libgomp.fortran/task-detach-4.f90 b/libgomp/testsuite/libgomp.fortran/task-detach-4.f90
index 6d0843c..159624c 100644
--- a/libgomp/testsuite/libgomp.fortran/task-detach-4.f90
+++ b/libgomp/testsuite/libgomp.fortran/task-detach-4.f90
@@ -11,9 +11,9 @@ program task_detach_4
 
   !$omp parallel
     !$omp single
-      !$omp task detach(detach_event)
+      !$omp task detach (detach_event)
         x = x + 1
-	call omp_fulfill_event(detach_event)
+	call omp_fulfill_event (detach_event)
       !$omp end task
     !$omp end single
   !$omp end parallel
diff --git a/libgomp/testsuite/libgomp.fortran/task-detach-5.f90 b/libgomp/testsuite/libgomp.fortran/task-detach-5.f90
index 955d687..95bd132 100644
--- a/libgomp/testsuite/libgomp.fortran/task-detach-5.f90
+++ b/libgomp/testsuite/libgomp.fortran/task-detach-5.f90
@@ -10,17 +10,17 @@ program task_detach_5
   integer :: x = 0, y = 0, z = 0
   integer :: thread_count
 
-  !$omp parallel firstprivate(detach_event1, detach_event2)
+  !$omp parallel private (detach_event1, detach_event2)
     !$omp single
-      thread_count = omp_get_num_threads()
+      thread_count = omp_get_num_threads ()
     !$omp end single
 
-    !$omp task detach(detach_event1) untied
+    !$omp task detach (detach_event1) untied
       !$omp atomic update
 	x = x + 1
     !$omp end task
 
-    !$omp task detach(detach_event2) untied
+    !$omp task detach (detach_event2) untied
       !$omp atomic update
 	y = y + 1
       call omp_fulfill_event (detach_event1);
diff --git a/libgomp/testsuite/libgomp.fortran/task-detach-6.f90 b/libgomp/testsuite/libgomp.fortran/task-detach-6.f90
index 0fe2155..b2c476f 100644
--- a/libgomp/testsuite/libgomp.fortran/task-detach-6.f90
+++ b/libgomp/testsuite/libgomp.fortran/task-detach-6.f90
@@ -11,30 +11,28 @@ program task_detach_6
   integer :: x = 0, y = 0, z = 0
   integer :: thread_count
 
-  !$omp target map(tofrom: x, y, z) map(from: thread_count)
-    !$omp parallel firstprivate(detach_event1, detach_event2)
+  !$omp target map (tofrom: x, y, z) map (from: thread_count)
+    !$omp parallel private (detach_event1, detach_event2)
       !$omp single
-	thread_count = omp_get_num_threads()
+	thread_count = omp_get_num_threads ()
       !$omp end single
 
-      !$omp task detach(detach_event1) untied
+      !$omp task detach (detach_event1) untied
 	!$omp atomic update
 	  x = x + 1
       !$omp end task
 
-      !$omp task detach(detach_event2) untied
+      !$omp task detach (detach_event2) untied
 	!$omp atomic update
 	  y = y + 1
-	call omp_fulfill_event (detach_event1);
+	call omp_fulfill_event (detach_event1)
       !$omp end task
 
       !$omp task untied
 	!$omp atomic update
 	  z = z + 1
-	call omp_fulfill_event (detach_event2);
+	call omp_fulfill_event (detach_event2)
       !$omp end task
-
-      !$omp taskwait
     !$omp end parallel
   !$omp end target
 
diff --git a/libgomp/testsuite/libgomp.fortran/task-detach-7.f90 b/libgomp/testsuite/libgomp.fortran/task-detach-7.f90
new file mode 100644
index 0000000..32e715e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/task-detach-7.f90
@@ -0,0 +1,42 @@
+! { dg-do run }
+
+! Test tasks with detach clause.  Each thread spawns off a chain of tasks,
+! that can then be executed by any available thread.  Each thread uses
+! taskwait to wait for the child tasks to complete.
+
+program task_detach_7
+  use omp_lib
+
+  integer (kind=omp_event_handle_kind) :: detach_event1, detach_event2
+  integer :: x = 0, y = 0, z = 0
+  integer :: thread_count
+
+  !$omp parallel private (detach_event1, detach_event2)
+    !$omp single
+      thread_count = omp_get_num_threads()
+    !$omp end single
+
+    !$omp task detach (detach_event1) untied
+      !$omp atomic update
+	x = x + 1
+    !$omp end task
+
+    !$omp task detach (detach_event2) untied
+      !$omp atomic update
+	y = y + 1
+      call omp_fulfill_event (detach_event1)
+    !$omp end task
+
+    !$omp task untied
+      !$omp atomic update
+	z = z + 1
+      call omp_fulfill_event (detach_event2)
+    !$omp end task
+
+    !$omp taskwait
+  !$omp end parallel
+
+  if (x /= thread_count) stop 1
+  if (y /= thread_count) stop 2
+  if (z /= thread_count) stop 3
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/task-detach-8.f90 b/libgomp/testsuite/libgomp.fortran/task-detach-8.f90
new file mode 100644
index 0000000..e760eab
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/task-detach-8.f90
@@ -0,0 +1,45 @@
+! { dg-do run }
+
+! 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.  Each thread uses taskwait to wait for the
+! child tasks to complete.
+
+program task_detach_8
+  use omp_lib
+
+  integer (kind=omp_event_handle_kind) :: detach_event1, detach_event2
+  integer :: x = 0, y = 0, z = 0
+  integer :: thread_count
+
+  !$omp target map (tofrom: x, y, z) map (from: thread_count)
+    !$omp parallel private (detach_event1, detach_event2)
+      !$omp single
+	thread_count = omp_get_num_threads ()
+      !$omp end single
+
+      !$omp task detach (detach_event1) untied
+	!$omp atomic update
+	  x = x + 1
+      !$omp end task
+
+      !$omp task detach (detach_event2) untied
+	!$omp atomic update
+	  y = y + 1
+	call omp_fulfill_event (detach_event1)
+      !$omp end task
+
+      !$omp task untied
+	!$omp atomic update
+	  z = z + 1
+	call omp_fulfill_event (detach_event2)
+      !$omp end task
+
+      !$omp taskwait
+    !$omp end parallel
+  !$omp end target
+
+  if (x /= thread_count) stop 1
+  if (y /= thread_count) stop 2
+  if (z /= thread_count) stop 3
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/task-detach-9.f90 b/libgomp/testsuite/libgomp.fortran/task-detach-9.f90
new file mode 100644
index 0000000..540c6de
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/task-detach-9.f90
@@ -0,0 +1,41 @@
+! { dg-do run }
+
+! Test tasks with detach clause.  Each thread spawns off a chain of tasks
+! in a taskgroup, that can then be executed by any available thread.
+
+program task_detach_9
+  use omp_lib
+
+  integer (kind=omp_event_handle_kind) :: detach_event1, detach_event2
+  integer :: x = 0, y = 0, z = 0
+  integer :: thread_count
+
+  !$omp parallel private (detach_event1, detach_event2)
+    !$omp taskgroup
+      !$omp single
+	thread_count = omp_get_num_threads ()
+      !$omp end single
+
+      !$omp task detach (detach_event1) untied
+	!$omp atomic update
+	  x = x + 1
+      !$omp end task
+
+      !$omp task detach (detach_event2) untied
+	!$omp atomic update
+	  y = y + 1
+	call omp_fulfill_event (detach_event1);
+      !$omp end task
+
+      !$omp task untied
+	!$omp atomic update
+	  z = z + 1
+	call omp_fulfill_event (detach_event2);
+      !$omp end task
+    !$omp end taskgroup
+  !$omp end parallel
+
+  if (x /= thread_count) stop 1
+  if (y /= thread_count) stop 2
+  if (z /= thread_count) stop 3
+end program