openmp: Add support for the OpenMP 5.0 task detach clause

2021-01-16  Kwok Cheung Yeung  <kcy@codesourcery.com>

	gcc/
	* builtin-types.def
	(BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT): Rename
	to...
	(BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT_PTR):
	...this.  Add extra argument.
	* gimplify.c (omp_default_clause): Ensure that event handle is
	firstprivate in a task region.
	(gimplify_scan_omp_clauses): Handle OMP_CLAUSE_DETACH.
	(gimplify_adjust_omp_clauses): Likewise.
	* omp-builtins.def (BUILT_IN_GOMP_TASK): Change function type to
	BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT_PTR.
	* omp-expand.c (expand_task_call): Add GOMP_TASK_FLAG_DETACH to flags
	if detach clause specified.  Add detach argument when generating
	call to	GOMP_task.
	* omp-low.c (scan_sharing_clauses): Setup data environment for detach
	clause.
	(finish_taskreg_scan): Move field for variable containing the event
	handle to the front of the struct.
	* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_DETACH.  Fix
	ordering.
	* tree-nested.c (convert_nonlocal_omp_clauses): Handle
	OMP_CLAUSE_DETACH clause.
	(convert_local_omp_clauses): Handle OMP_CLAUSE_DETACH clause.
	* tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE_DETACH.
	* tree.c (omp_clause_num_ops): Add entry for OMP_CLAUSE_DETACH.
	Fix ordering.
	(omp_clause_code_name): Add entry for OMP_CLAUSE_DETACH.  Fix
	ordering.
	(walk_tree_1): Handle OMP_CLAUSE_DETACH.

	gcc/c-family/
	* c-pragma.h (pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_DETACH.
	Redefine PRAGMA_OACC_CLAUSE_DETACH.

	gcc/c/
	* c-parser.c (c_parser_omp_clause_detach): New.
	(c_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_DETACH clause.
	(OMP_TASK_CLAUSE_MASK): Add mask for PRAGMA_OMP_CLAUSE_DETACH.
	* c-typeck.c (c_finish_omp_clauses): Handle PRAGMA_OMP_CLAUSE_DETACH
	clause.  Prevent use of detach with mergeable and overriding the
	data sharing mode of the event handle.

	gcc/cp/
	* parser.c (cp_parser_omp_clause_detach): New.
	(cp_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_DETACH.
	(OMP_TASK_CLAUSE_MASK): Add mask for PRAGMA_OMP_CLAUSE_DETACH.
	* pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_DETACH clause.
	* semantics.c (finish_omp_clauses): Handle OMP_CLAUSE_DETACH clause.
	Prevent use of detach with mergeable and overriding the	data sharing
	mode of the event handle.

	gcc/fortran/
	* dump-parse-tree.c (show_omp_clauses): Handle detach clause.
	* frontend-passes.c (gfc_code_walker): Walk detach expression.
	* gfortran.h (struct gfc_omp_clauses): Add detach field.
	(gfc_c_intptr_kind): New.
	* openmp.c (gfc_free_omp_clauses): Free detach clause.
	(gfc_match_omp_detach): New.
	(enum omp_mask1): Add OMP_CLAUSE_DETACH.
	(enum omp_mask2): Remove OMP_CLAUSE_DETACH.
	(gfc_match_omp_clauses): Handle OMP_CLAUSE_DETACH for OpenMP.
	(OMP_TASK_CLAUSES): Add OMP_CLAUSE_DETACH.
	(resolve_omp_clauses): Prevent use of detach with mergeable and
	overriding the data sharing mode of the event handle.
	* trans-openmp.c (gfc_trans_omp_clauses): Handle detach clause.
	* trans-types.c (gfc_c_intptr_kind): New.
	(gfc_init_kinds): Initialize gfc_c_intptr_kind.
	* types.def
	(BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT): Rename
	to...
	(BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT_PTR):
	...this.  Add extra argument.

	gcc/testsuite/
	* c-c++-common/gomp/task-detach-1.c: New.
	* g++.dg/gomp/task-detach-1.C: New.
	* gcc.dg/gomp/task-detach-1.c: New.
	* gfortran.dg/gomp/task-detach-1.f90: New.

	include/
	* gomp-constants.h (GOMP_TASK_FLAG_DETACH): New.

	libgomp/
	* fortran.c (omp_fulfill_event_): New.
	* libgomp.h (struct gomp_task): Add detach and completion_sem fields.
	(struct gomp_team): Add task_detach_queue and task_detach_count
	fields.
	* libgomp.map (OMP_5.0.1): Add omp_fulfill_event and omp_fulfill_event_.
	* libgomp_g.h (GOMP_task): Add extra argument.
	* omp.h.in (enum omp_event_handle_t): New.
	(omp_fulfill_event): New.
	* omp_lib.f90.in (omp_event_handle_kind): New.
	(omp_fulfill_event): New.
	* omp_lib.h.in (omp_event_handle_kind): New.
	(omp_fulfill_event): Declare.
	* priority_queue.c (priority_tree_find): New.
	(priority_list_find): New.
	(priority_queue_find): New.
	* priority_queue.h (priority_queue_predicate): New.
	(priority_queue_find): New.
	* task.c (gomp_init_task): Initialize detach field.
	(task_fulfilled_p): New.
	(GOMP_task): Add detach argument.  Ignore detach argument if
	GOMP_TASK_FLAG_DETACH not set in flags.  Initialize completion_sem
	field.	Copy address of completion_sem into detach argument and
	into the start of the data record.  Wait for detach event if task
	not deferred.
	(gomp_barrier_handle_tasks): Queue tasks with unfulfilled events.
	Remove completed tasks and requeue dependent tasks.
	(omp_fulfill_event): New.
	* team.c (gomp_new_team): Initialize task_detach_queue and
	task_detach_count fields.
	(free_team): Free task_detach_queue field.
	* testsuite/libgomp.c-c++-common/task-detach-1.c: New testcase.
	* testsuite/libgomp.c-c++-common/task-detach-2.c: New testcase.
	* testsuite/libgomp.c-c++-common/task-detach-3.c: New testcase.
	* testsuite/libgomp.c-c++-common/task-detach-4.c: New testcase.
	* testsuite/libgomp.c-c++-common/task-detach-5.c: New testcase.
	* testsuite/libgomp.c-c++-common/task-detach-6.c: New testcase.
	* testsuite/libgomp.fortran/task-detach-1.f90: New testcase.
	* testsuite/libgomp.fortran/task-detach-2.f90: New testcase.
	* testsuite/libgomp.fortran/task-detach-3.f90: New testcase.
	* testsuite/libgomp.fortran/task-detach-4.f90: New testcase.
	* testsuite/libgomp.fortran/task-detach-5.f90: New testcase.
	* testsuite/libgomp.fortran/task-detach-6.f90: New testcase.

23 files changed, 636 insertions, 22 deletions

diff --git a/libgomp/fortran.c b/libgomp/fortran.c
index e6bccb9..4ec39c4 100644
--- a/libgomp/fortran.c
+++ b/libgomp/fortran.c
@@ -605,6 +605,12 @@ omp_get_max_task_priority_ (void)
 }
 
 void
+omp_fulfill_event_ (intptr_t event)
+{
+  omp_fulfill_event ((omp_event_handle_t) event);
+}
+
+void
 omp_set_affinity_format_ (const char *format, size_t format_len)
 {
   gomp_set_affinity_format (format, format_len);
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 305cba3..b4d0c93 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -545,6 +545,9 @@ 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;
@@ -685,6 +688,10 @@ 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;
+  unsigned int task_detach_count;
+
   /* This array contains structures for implicit tasks.  */
   struct gomp_task implicit_task[];
 };
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index 2c95f78..4ad190a 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -195,6 +195,8 @@ OMP_5.0.1 {
 	omp_free;
 	omp_get_supported_active_levels;
 	omp_get_supported_active_levels_;
+	omp_fulfill_event;
+	omp_fulfill_event_;
 } OMP_5.0;
 
 GOMP_1.0 {
diff --git a/libgomp/libgomp_g.h b/libgomp/libgomp_g.h
index 0e1fbee..3cbe0a4 100644
--- a/libgomp/libgomp_g.h
+++ b/libgomp/libgomp_g.h
@@ -294,7 +294,7 @@ extern bool GOMP_cancellation_point (int);
 /* task.c */
 
 extern void GOMP_task (void (*) (void *), void *, void (*) (void *, void *),
-		       long, long, bool, unsigned, void **, int);
+		       long, long, bool, unsigned, void **, int, void *);
 extern void GOMP_taskloop (void (*) (void *), void *,
 			   void (*) (void *, void *), long, long, unsigned,
 			   unsigned long, int, long, long, long);
diff --git a/libgomp/omp.h.in b/libgomp/omp.h.in
index f88e360..69f96f0 100644
--- a/libgomp/omp.h.in
+++ b/libgomp/omp.h.in
@@ -171,6 +171,11 @@ typedef struct omp_alloctrait_t
   omp_uintptr_t value;
 } omp_alloctrait_t;
 
+typedef enum omp_event_handle_t __GOMP_UINTPTR_T_ENUM
+{
+  __omp_event_handle_t_max__ = __UINTPTR_MAX__
+} omp_event_handle_t;
+
 #ifdef __cplusplus
 extern "C" {
 # define __GOMP_NOTHROW throw ()
@@ -245,6 +250,8 @@ extern int omp_is_initial_device (void) __GOMP_NOTHROW;
 extern int omp_get_initial_device (void) __GOMP_NOTHROW;
 extern int omp_get_max_task_priority (void) __GOMP_NOTHROW;
 
+extern void omp_fulfill_event (omp_event_handle_t) __GOMP_NOTHROW;
+
 extern void *omp_target_alloc (__SIZE_TYPE__, int) __GOMP_NOTHROW;
 extern void omp_target_free (void *, int) __GOMP_NOTHROW;
 extern int omp_target_is_present (const void *, int) __GOMP_NOTHROW;
diff --git a/libgomp/omp_lib.f90.in b/libgomp/omp_lib.f90.in
index ff00afa..851f85f 100644
--- a/libgomp/omp_lib.f90.in
+++ b/libgomp/omp_lib.f90.in
@@ -39,6 +39,7 @@
         integer, parameter :: omp_alloctrait_val_kind = c_intptr_t
         integer, parameter :: omp_memspace_handle_kind = c_intptr_t
         integer, parameter :: omp_depend_kind = @OMP_DEPEND_KIND@
+        integer, parameter :: omp_event_handle_kind = c_intptr_t
         integer (omp_sched_kind), parameter :: omp_sched_static = 1
         integer (omp_sched_kind), parameter :: omp_sched_dynamic = 2
         integer (omp_sched_kind), parameter :: omp_sched_guided = 3
@@ -556,6 +557,14 @@
         end interface
 
         interface
+          subroutine omp_fulfill_event (event)
+            use omp_lib_kinds
+            integer (kind=omp_event_handle_kind), &
+              value, intent(in) :: event
+          end subroutine omp_fulfill_event
+        end interface
+
+        interface
           subroutine omp_set_affinity_format (format)
             character(len=*), intent(in) :: format
           end subroutine omp_set_affinity_format
diff --git a/libgomp/omp_lib.h.in b/libgomp/omp_lib.h.in
index a00d9bd..06d17b5 100644
--- a/libgomp/omp_lib.h.in
+++ b/libgomp/omp_lib.h.in
@@ -82,10 +82,12 @@
 
       integer omp_allocator_handle_kind, omp_alloctrait_key_kind
       integer omp_alloctrait_val_kind, omp_memspace_handle_kind
+      integer omp_event_handle_kind
       parameter (omp_allocator_handle_kind = @INTPTR_T_KIND@)
       parameter (omp_alloctrait_key_kind = 4)
       parameter (omp_alloctrait_val_kind = @INTPTR_T_KIND@)
       parameter (omp_memspace_handle_kind = @INTPTR_T_KIND@)
+      parameter (omp_event_handle_kind = @INTPTR_T_KIND@)
       integer (omp_alloctrait_key_kind) omp_atk_sync_hint
       integer (omp_alloctrait_key_kind) omp_atk_alignment
       integer (omp_alloctrait_key_kind) omp_atk_access
@@ -245,6 +247,8 @@
       external omp_get_max_task_priority
       integer(4) omp_get_max_task_priority
 
+      external omp_fulfill_event
+
       external omp_set_affinity_format, omp_get_affinity_format
       external omp_display_affinity, omp_capture_affinity
       integer(4) omp_get_affinity_format
diff --git a/libgomp/priority_queue.c b/libgomp/priority_queue.c
index 6361f36..39b69f4 100644
--- a/libgomp/priority_queue.c
+++ b/libgomp/priority_queue.c
@@ -168,6 +168,63 @@ priority_queue_verify (enum priority_queue_type type,
 }
 #endif /* _LIBGOMP_CHECKING_ */
 
+/* Tree version of priority_queue_find.  */
+
+static struct gomp_task *
+priority_tree_find (enum priority_queue_type type,
+		    prio_splay_tree_node node,
+		    priority_queue_predicate pred)
+{
+ again:
+  if (!node)
+    return NULL;
+  struct gomp_task *task = priority_tree_find (type, node->right, pred);
+  if (task)
+    return task;
+  task = priority_node_to_task (type, node->key.l.tasks);
+  if (pred (task))
+    return task;
+  node = node->left;
+  goto again;
+}
+
+/* List version of priority_queue_find.  */
+
+static struct gomp_task *
+priority_list_find (enum priority_queue_type type,
+		     struct priority_list *list,
+		     priority_queue_predicate pred)
+{
+  struct priority_node *node = list->tasks;
+  if (!node)
+    return NULL;
+
+  do
+    {
+      struct gomp_task *task = priority_node_to_task (type, node);
+      if (pred (task))
+	return task;
+      node = node->next;
+    }
+  while (node != list->tasks);
+
+  return NULL;
+}
+
+/* Return the highest priority task in the priority queue HEAD that
+   satisfies the predicate PRED.  HEAD contains tasks of type TYPE.  */
+
+struct gomp_task *
+priority_queue_find (enum priority_queue_type type,
+		     struct priority_queue *head,
+		     priority_queue_predicate pred)
+{
+  if (priority_queue_multi_p (head))
+    return priority_tree_find (type, head->t.root, pred);
+  else
+    return priority_list_find (type, &head->l, pred);
+}
+
 /* Remove NODE from priority queue HEAD, wherever it may be inside the
    tree.  HEAD contains tasks of type TYPE.  */
 
diff --git a/libgomp/priority_queue.h b/libgomp/priority_queue.h
index 41f5c73..d8d31b7 100644
--- a/libgomp/priority_queue.h
+++ b/libgomp/priority_queue.h
@@ -113,6 +113,8 @@ enum priority_queue_type
   PQ_IGNORED = 999
 };
 
+typedef bool (*priority_queue_predicate) (struct gomp_task *);
+
 /* Priority queue implementation prototypes.  */
 
 extern bool priority_queue_task_in_queue_p (enum priority_queue_type,
@@ -122,6 +124,9 @@ extern void priority_queue_dump (enum priority_queue_type,
 				 struct priority_queue *);
 extern void priority_queue_verify (enum priority_queue_type,
 				   struct priority_queue *, bool);
+extern struct gomp_task *priority_queue_find (enum priority_queue_type,
+					      struct priority_queue *,
+					      priority_queue_predicate);
 extern void priority_tree_remove (enum priority_queue_type,
 				  struct priority_queue *,
 				  struct priority_node *);
diff --git a/libgomp/task.c b/libgomp/task.c
index 0e9887d..5ece878 100644
--- a/libgomp/task.c
+++ b/libgomp/task.c
@@ -86,6 +86,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;
 }
 
 /* Clean up a task, after completing it.  */
@@ -326,6 +327,12 @@ gomp_task_handle_depend (struct gomp_task *task, struct gomp_task *parent,
     }
 }
 
+static bool
+task_fulfilled_p (struct gomp_task *task)
+{
+  return __atomic_load_n (&task->completion_sem, __ATOMIC_RELAXED);
+}
+
 /* 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.
@@ -347,7 +354,7 @@ gomp_task_handle_depend (struct gomp_task *task, struct gomp_task *parent,
 void
 GOMP_task (void (*fn) (void *), void *data, void (*cpyfn) (void *, void *),
 	   long arg_size, long arg_align, bool if_clause, unsigned flags,
-	   void **depend, int priority)
+	   void **depend, int priority, void *detach)
 {
   struct gomp_thread *thr = gomp_thread ();
   struct gomp_team *team = thr->ts.team;
@@ -383,6 +390,9 @@ GOMP_task (void (*fn) (void *), void *data, void (*cpyfn) (void *, void *),
   else if (priority > gomp_max_task_priority_var)
     priority = gomp_max_task_priority_var;
 
+  if ((flags & GOMP_TASK_FLAG_DETACH) == 0)
+    detach = NULL;
+
   if (!if_clause || team == NULL
       || (thr->task && thr->task->final_task)
       || team->task_count > 64 * team->nthreads)
@@ -404,6 +414,18 @@ GOMP_task (void (*fn) (void *), void *data, void (*cpyfn) (void *, void *),
       task.final_task = (thr->task && thr->task->final_task)
 			|| (flags & GOMP_TASK_FLAG_FINAL);
       task.priority = priority;
+
+      if (detach)
+	{
+	  task.detach = true;
+	  gomp_sem_init (&task.completion_sem, 0);
+	  *(void **) detach = &task.completion_sem;
+	  if (data)
+	    *(void **) data = &task.completion_sem;
+
+	  gomp_debug (0, "New event: %p\n", &task.completion_sem);
+	}
+
       if (thr->task)
 	{
 	  task.in_tied_task = thr->task->in_tied_task;
@@ -420,6 +442,10 @@ GOMP_task (void (*fn) (void *), void *data, void (*cpyfn) (void *, void *),
 	}
       else
 	fn (data);
+
+      if (detach && !task_fulfilled_p (&task))
+	gomp_sem_wait (&task.completion_sem);
+
       /* Access to "children" is normally done inside a task_lock
 	 mutex region, but the only way this particular task.children
 	 can be set is if this thread's task work function (fn)
@@ -458,6 +484,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;
+      if (detach)
+	{
+	  task->detach = true;
+	  gomp_sem_init (&task->completion_sem, 0);
+	  *(void **) detach = &task->completion_sem;
+	  if (data)
+	    *(void **) data = &task->completion_sem;
+
+	  gomp_debug (0, "New event: %p\n", &task->completion_sem);
+	}
       thr->task = task;
       if (cpyfn)
 	{
@@ -1325,6 +1361,28 @@ gomp_barrier_handle_tasks (gomp_barrier_state_t state)
   while (1)
     {
       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;
@@ -1392,29 +1450,43 @@ gomp_barrier_handle_tasks (gomp_barrier_state_t state)
       gomp_mutex_lock (&team->task_lock);
       if (child_task)
 	{
-	 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)
+	  if (child_task->detach && !task_fulfilled_p (child_task))
 	    {
-	      do_wake = team->nthreads - team->task_running_count;
-	      if (do_wake > new_tasks)
-		do_wake = new_tasks;
+	      priority_queue_insert (PQ_TEAM, &team->task_detach_queue,
+				     child_task, child_task->priority,
+				     PRIORITY_INSERT_END,
+				     false, false);
+	      ++team->task_detach_count;
+	      gomp_debug (0, "thread %d: queueing task with event %p\n",
+			  thr->ts.team_id, &child_task->completion_sem);
+	      child_task = NULL;
 	    }
-	  if (--team->task_count == 0
-	      && gomp_team_barrier_waiting_for_tasks (&team->barrier))
+	  else
 	    {
-	      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);
+	     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);
+		}
 	    }
 	}
     }
@@ -2326,3 +2398,21 @@ omp_in_final (void)
 }
 
 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;
+
+  if (__atomic_load_n (sem, __ATOMIC_RELAXED))
+    gomp_fatal ("omp_fulfill_event: %p event already fulfilled!\n", sem);
+
+  gomp_debug (0, "omp_fulfill_event: %p\n", sem);
+  gomp_sem_post (sem);
+  if (team)
+    gomp_team_barrier_wake (&team->barrier, 1);
+}
+
+ialias (omp_fulfill_event)
diff --git a/libgomp/team.c b/libgomp/team.c
index 85d5305..0f3707c 100644
--- a/libgomp/team.c
+++ b/libgomp/team.c
@@ -206,6 +206,9 @@ 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;
 }
 
@@ -221,6 +224,7 @@ 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
new file mode 100644
index 0000000..8583e37
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-1.c
@@ -0,0 +1,36 @@
+/* { dg-do run } */
+
+#include <omp.h>
+#include <assert.h>
+
+/* Test chaining of detached tasks, with each task fulfilling the
+   completion event of the previous one.  */
+
+int main (void)
+{
+  omp_event_handle_t detach_event1, detach_event2;
+  int x = 0, y = 0, z = 0;
+
+  #pragma omp parallel
+    #pragma omp single
+    {
+      #pragma omp task detach(detach_event1)
+	x++;
+
+      #pragma omp task detach(detach_event2)
+      {
+	y++;
+	omp_fulfill_event (detach_event1);
+      }
+
+      #pragma omp task
+      {
+	z++;
+	omp_fulfill_event (detach_event2);
+      }
+    }
+
+  assert (x == 1);
+  assert (y == 1);
+  assert (z == 1);
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/task-detach-2.c b/libgomp/testsuite/libgomp.c-c++-common/task-detach-2.c
new file mode 100644
index 0000000..943ac2a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-2.c
@@ -0,0 +1,37 @@
+/* { dg-do run } */
+
+#include <omp.h>
+#include <assert.h>
+
+/* Test handling of detach clause with only a single thread.  The runtime
+   should not block when a task with an unfulfilled event finishes
+   running.  */
+
+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 single
+    {
+      #pragma omp task detach(detach_event1)
+	x++;
+
+      #pragma omp task detach(detach_event2)
+      {
+	y++;
+	omp_fulfill_event (detach_event1);
+      }
+
+      #pragma omp task
+      {
+	z++;
+	omp_fulfill_event (detach_event2);
+      }
+    }
+
+  assert (x == 1);
+  assert (y == 1);
+  assert (z == 1);
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/task-detach-3.c b/libgomp/testsuite/libgomp.c-c++-common/task-detach-3.c
new file mode 100644
index 0000000..2609fb1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-3.c
@@ -0,0 +1,33 @@
+/* { dg-do run } */
+
+#include <omp.h>
+#include <assert.h>
+
+/* Test the task detach clause used together with dependencies.  */
+
+int main (void)
+{
+  omp_event_handle_t detach_event;
+  int x = 0, y = 0, z = 0;
+  int dep;
+
+  #pragma omp parallel
+    #pragma omp single
+    {
+      #pragma omp task depend(out:dep) detach(detach_event)
+	x++;
+
+      #pragma omp task
+      {
+	y++;
+	omp_fulfill_event(detach_event);
+      }
+
+      #pragma omp task depend(in:dep)
+	z++;
+    }
+
+  assert (x == 1);
+  assert (y == 1);
+  assert (z == 1);
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/task-detach-4.c b/libgomp/testsuite/libgomp.c-c++-common/task-detach-4.c
new file mode 100644
index 0000000..eeb9554
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-4.c
@@ -0,0 +1,24 @@
+/* { dg-do run } */
+
+#include <omp.h>
+#include <assert.h>
+
+/* Test detach clause, where a task fulfills its own completion event.  */
+
+int main (void)
+{
+  omp_event_handle_t detach_event;
+  int x = 0;
+
+  detach_event = (omp_event_handle_t) 0x123456789abcdef0;
+
+  #pragma omp parallel
+    #pragma omp single
+      #pragma omp task detach(detach_event)
+      {
+	x++;
+	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
new file mode 100644
index 0000000..5a01517
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-5.c
@@ -0,0 +1,42 @@
+/* { 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.  */
+
+int main (void)
+{
+  int x = 0, y = 0, z = 0;
+  int thread_count;
+  omp_event_handle_t detach_event1, detach_event2;
+
+  #pragma omp parallel firstprivate(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);
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/task-detach-6.c b/libgomp/testsuite/libgomp.c-c++-common/task-detach-6.c
new file mode 100644
index 0000000..b5f68cc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-6.c
@@ -0,0 +1,46 @@
+/* { 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 firstprivate(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.fortran/task-detach-1.f90 b/libgomp/testsuite/libgomp.fortran/task-detach-1.f90
new file mode 100644
index 0000000..217bf65
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/task-detach-1.f90
@@ -0,0 +1,33 @@
+! { dg-do run }
+
+! Test chaining of detached tasks, with each task fulfilling the
+! completion event of the previous one.
+
+program task_detach_1
+  use omp_lib
+
+  integer (kind=omp_event_handle_kind) :: detach_event1, detach_event2
+  integer :: x = 0, y = 0, z = 0
+
+  !$omp parallel
+    !$omp single
+      !$omp task detach(detach_event1)
+        x = x + 1
+      !$omp end task
+
+      !$omp task detach(detach_event2)
+        y = y + 1
+	call omp_fulfill_event (detach_event1)
+      !$omp end task
+
+      !$omp task
+        z = z + 1
+	call omp_fulfill_event (detach_event2)
+      !$omp end task
+    !$omp end single
+  !$omp end parallel
+
+  if (x /= 1) stop 1
+  if (y /= 1) stop 2
+  if (z /= 1) stop 3
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/task-detach-2.f90 b/libgomp/testsuite/libgomp.fortran/task-detach-2.f90
new file mode 100644
index 0000000..ecb4829
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/task-detach-2.f90
@@ -0,0 +1,34 @@
+! { dg-do run }
+
+! Test handling of detach clause with only a single thread.  The runtime
+! should not block when a task with an unfulfilled event finishes
+! running.
+
+program task_detach_2
+  use omp_lib
+
+  integer (kind=omp_event_handle_kind) :: detach_event1, detach_event2
+  integer :: x = 0, y = 0, z = 0
+
+  !$omp parallel num_threads(1)
+    !$omp single
+      !$omp task detach(detach_event1)
+        x = x + 1
+      !$omp end task
+
+      !$omp task detach(detach_event2)
+        y = y + 1
+	call omp_fulfill_event (detach_event1)
+      !$omp end task
+
+      !$omp task
+        z = z + 1
+	call omp_fulfill_event (detach_event2)
+      !$omp end task
+    !$omp end single
+  !$omp end parallel
+
+  if (x /= 1) stop 1
+  if (y /= 1) stop 2
+  if (z /= 1) stop 3
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/task-detach-3.f90 b/libgomp/testsuite/libgomp.fortran/task-detach-3.f90
new file mode 100644
index 0000000..bdf93a5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/task-detach-3.f90
@@ -0,0 +1,33 @@
+! { dg-do run }
+
+! Test the task detach clause used together with dependencies.
+
+program task_detach_3
+
+  use omp_lib
+
+  integer (kind=omp_event_handle_kind) :: detach_event
+  integer :: x = 0, y = 0, z = 0
+  integer :: dep
+
+  !$omp parallel
+    !$omp single
+      !$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)
+      !$omp end task
+
+      !$omp task depend(in:dep)
+        z = z + 1
+      !$omp end task
+    !$omp end single
+  !$omp end parallel
+
+  if (x /= 1) stop 1
+  if (y /= 1) stop 2
+  if (z /= 1) stop 3
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/task-detach-4.f90 b/libgomp/testsuite/libgomp.fortran/task-detach-4.f90
new file mode 100644
index 0000000..6d0843c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/task-detach-4.f90
@@ -0,0 +1,22 @@
+! { dg-do run }
+
+! Test detach clause, where a task fulfills its own completion event.
+
+program task_detach_4
+
+  use omp_lib
+
+  integer (kind=omp_event_handle_kind) :: detach_event
+  integer :: x = 0
+
+  !$omp parallel
+    !$omp single
+      !$omp task detach(detach_event)
+        x = x + 1
+	call omp_fulfill_event(detach_event)
+      !$omp end task
+    !$omp end single
+  !$omp end parallel
+
+  if (x /= 1) stop 1
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/task-detach-5.f90 b/libgomp/testsuite/libgomp.fortran/task-detach-5.f90
new file mode 100644
index 0000000..955d687
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/task-detach-5.f90
@@ -0,0 +1,39 @@
+! { 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.
+
+program task_detach_5
+  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 firstprivate(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 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-6.f90 b/libgomp/testsuite/libgomp.fortran/task-detach-6.f90
new file mode 100644
index 0000000..0fe2155
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/task-detach-6.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, that can then be executed by
+! any available thread.
+
+program task_detach_6
+  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 firstprivate(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