diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index b4d0c93dd20..ef1bb4907b6 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
@@ -537,6 +540,16 @@ struct gomp_task
      into the various queues to be scheduled.  */
   size_t num_dependees;
 
+  union {
+      /* Valid only if deferred_p is false.  */
+      gomp_sem_t *completion_sem;
+      /* Valid only if deferred_p is true.  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;
+    };
+  bool deferred_p;
+
   /* Priority of this task.  */
   int priority;
   /* The priority node for this task in each of the different queues.
@@ -545,9 +558,6 @@ struct gomp_task
      entries and the gomp_task in which they reside.  */
   struct priority_node pnode[3];
 
-  bool detach;
-  gomp_sem_t completion_sem;
-
   struct gomp_task_icv icv;
   void (*fn) (void *);
   void *fn_data;
@@ -688,8 +698,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 b242e7c8d20..1c73c759a8d 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;
@@ -74,19 +75,20 @@ gomp_init_task (struct gomp_task *task, struct gomp_task *parent_task,
      benchmark the overhead of creating tasks as there are millions of
      tiny tasks created that all run undeferred.  */
   task->parent = parent_task;
-  task->icv = *prev_icv;
-  task->kind = GOMP_TASK_IMPLICIT;
-  task->taskwait = NULL;
-  task->in_tied_task = false;
-  task->final_task = false;
-  task->copy_ctors_done = false;
-  task->parent_depends_on = false;
   priority_queue_init (&task->children_queue);
   task->taskgroup = NULL;
   task->dependers = NULL;
   task->depend_hash = NULL;
+  task->taskwait = NULL;
   task->depend_count = 0;
-  task->detach = false;
+  task->completion_sem = NULL;
+  task->deferred_p = false;
+  task->icv = *prev_icv;
+  task->kind = GOMP_TASK_IMPLICIT;
+  task->in_tied_task = false;
+  task->final_task = false;
+  task->copy_ctors_done = false;
+  task->parent_depends_on = false;
 }
 
 /* Clean up a task, after completing it.  */
@@ -327,12 +329,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.
@@ -398,6 +394,7 @@ GOMP_task (void (*fn) (void *), void *data, void (*cpyfn) (void *, void *),
       || team->task_count > 64 * team->nthreads)
     {
       struct gomp_task task;
+      gomp_sem_t completion_sem;
 
       /* If there are depend clauses and earlier deferred sibling tasks
 	 with depend clauses, check if there isn't a dependency.  If there
@@ -417,13 +414,14 @@ 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;
+	  gomp_sem_init (&completion_sem, 0);
+	  task.completion_sem = &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,8 +441,11 @@ GOMP_task (void (*fn) (void *), void *data, void (*cpyfn) (void *, void *),
       else
 	fn (data);
 
-      if (task.detach && !task_fulfilled_p (&task))
-	gomp_sem_wait (&task.completion_sem);
+      if ((flags & GOMP_TASK_FLAG_DETACH) != 0)
+	{
+	  gomp_sem_wait (&completion_sem);
+	  gomp_sem_destroy (&completion_sem);
+	}
 
       /* Access to "children" is normally done inside a task_lock
 	 mutex region, but the only way this particular task.children
@@ -484,15 +485,16 @@ GOMP_task (void (*fn) (void *), void *data, void (*cpyfn) (void *, void *),
       task->kind = GOMP_TASK_UNDEFERRED;
       task->in_tied_task = parent->in_tied_task;
       task->taskgroup = taskgroup;
+      task->deferred_p = true;
       if ((flags & GOMP_TASK_FLAG_DETACH) != 0)
 	{
-	  task->detach = true;
-	  gomp_sem_init (&task->completion_sem, 0);
-	  *(void **) detach = &task->completion_sem;
-	  if (data)
-	    *(void **) data = &task->completion_sem;
+	  task->detach_team = team;
 
-	  gomp_debug (0, "New event: %p\n", &task->completion_sem);
+	  *(void **) detach = task;
+	  if (data)
+	    *(void **) data = task;
+
+	  gomp_debug (0, "Thread %d: new event: %p\n", thr->ts.team_id, task);
 	}
       thr->task = task;
       if (cpyfn)
@@ -1362,27 +1364,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;
@@ -1405,6 +1386,19 @@ gomp_barrier_handle_tasks (gomp_barrier_state_t state)
 	  team->task_running_count++;
 	  child_task->in_tied_task = true;
 	}
+      else 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);
+	  if (to_free)
+	    {
+	      gomp_finish_task (to_free);
+	      free (to_free);
+	    }
+	  return;
+	}
       gomp_mutex_unlock (&team->task_lock);
       if (do_wake)
 	{
@@ -1450,44 +1444,37 @@ 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;
 	    }
+	  --team->task_count;
 	}
     }
 }
@@ -1559,7 +1546,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 +1602,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 +2070,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 +2416,75 @@ 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->deferred_p)
+    {
+      if (gomp_sem_getcount (task->completion_sem) > 0)
+	gomp_fatal ("omp_fulfill_event: %p event already fulfilled!\n", task);
 
-  if (gomp_sem_getcount (sem) > 0)
-    gomp_fatal ("omp_fulfill_event: %p event already fulfilled!\n", sem);
+      gomp_debug (0, "omp_fulfill_event: %p event for undeferred task\n",
+		  task);
+      gomp_sem_post (task->completion_sem);
+      return;
+    }
 
-  gomp_debug (0, "omp_fulfill_event: %p\n", sem);
-  gomp_sem_post (sem);
-  if (team)
-    gomp_team_barrier_wake (&team->barrier, 1);
+  struct gomp_team *team = __atomic_load_n (&task->detach_team,
+					    MEMMODEL_RELAXED);
+  if (!team)
+    gomp_fatal ("omp_fulfill_event: %p event is invalid or has already "
+		"been fulfilled!\n", task);
+
+  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);
+      __atomic_store_n (&task->detach_team, NULL, MEMMODEL_RELAXED);
+      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--;
+
+  int do_wake = 0;
+  bool shackled_thread_p = team == gomp_thread ()->ts.team;
+  if (new_tasks > 0)
+    {
+      /* Wake up threads to run new tasks.  */
+      do_wake = team->nthreads - team->task_running_count;
+      if (do_wake > new_tasks)
+	do_wake = new_tasks;
+    }
+
+  if (!shackled_thread_p
+      && !do_wake
+      && team->task_detach_count == 0
+      && gomp_team_barrier_waiting_for_tasks (&team->barrier))
+    /* Ensure that at least one thread is woken up to signal that the
+       barrier can finish.  */
+    do_wake = 1;
+
+  /* If we are running in an unshackled thread, the team might vanish before
+     gomp_team_barrier_wake is run if we release the lock first, so keep the
+     lock for the call in that case.  */
+  if (shackled_thread_p)
+    gomp_mutex_unlock (&team->task_lock);
+  if (do_wake)
+    gomp_team_barrier_wake (&team->barrier, do_wake);
+  if (!shackled_thread_p)
+    gomp_mutex_unlock (&team->task_lock);
+
+  gomp_finish_task (task);
+  free (task);
 }
 
 ialias (omp_fulfill_event)
diff --git a/libgomp/team.c b/libgomp/team.c
index 0f3707c2f80..96622345b10 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 8583e3761b5..14932b02b9c 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 00000000000..10d6746f61a
--- /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 00000000000..dd002dcc812
--- /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 943ac2a3f85..3e33c409fe8 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 2609fb1c43d..c85857de21c 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 eeb9554c665..cd0d2b353d6 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 5a0151729ef..382f377881f 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 b5f68ccabdc..e5c2291e6ff 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 00000000000..3f025d6a13f
--- /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 00000000000..6f77f1263ab
--- /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 00000000000..5316ca545b5
--- /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 217bf657506..c53b1ca0029 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 00000000000..61f0ea8ba0a
--- /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 00000000000..b33baff7f68
--- /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 ecb4829e14b..68e3ff2ef81 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 bdf93a5909a..5ac68d55802 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 6d0843cb126..159624cc74b 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 955d687ca8b..95bd1325e41 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 0fe21553494..b2c476fd6a6 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 00000000000..32e715e0f81
--- /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 00000000000..e760eabe158
--- /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 00000000000..540c6deaef3
--- /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