Hello
This is a further update of the patch for task detach support.
- The memory for the event is not mapped on the target. This means that if
omp_fulfill_event is called from an 'omp target' section with a target that
does not share memory with the host, the event will not be fulfilled (and a
segfault will probably occur).
I was thinking of something along the lines of:
#pragma omp task detach (event)
{
}
#pragma omp target
{
omp_fulfill_event (event);
}
Would something like this be expected to work? I cannot find many examples of
the detach clause online, and none of them use any offloading constructs.
I have asked on the omp-lang mailing list - this is not expected to work.
- The tasks awaiting event fulfillment currently wait until there are no other
runnable tasks left. A better approach would be to poll (without blocking) the
waiting tasks whenever any task completes, immediately removing any
now-complete tasks and requeuing any dependent tasks.
This has now been implemented. On every iteration of the main loop in
gomp_barrier_handle_tasks, it first checks to see if any tasks in the detach
queue have a fulfilled completion event, and if so it will remove the task and
requeue any dependent tasks.
I have found another problem with the original blocking approach when the tasks
are on offload devices. On Nvidia and GCN, a bar.sync/s_barrier instruction is
issued when gomp_team_barrier_wake is called to synchronise the threads.
However, if some of the barrier threads are stuck waiting for semaphores
associated with completion events, and the fulfillment of those events are in
other tasks waiting to run, then the result is a deadlock as the threads cannot
synchronise without all the semaphores being released.
I have removed the blocking path on gomp_barrier_handle_tasks altogether, and
omp_fulfill_event now directly wakes the barrier threads to process any tasks
that are now complete.
I have also ensured that the event handle specified on the detach clause is
firstprivate by default on enclosing scopes.
I believe this patch is largely complete now. I have done a bootstrap on x86_64
and run the testsuites with no regressions. I have also run the libgomp
testsuite with offloading to Nvidia and AMD GCN devices, also with no
regressions. Is this patch okay for trunk (or would it be more appropriate to
wait until GCC 11 is branched off)?
Thanks
Kwok
commit 3d82db0fc3623e9dc241bed4c4cfd266574d45e7
Author: Kwok Cheung Yeung <k...@codesourcery.com>
Date: Wed Dec 9 09:33:46 2020 -0800
openmp: Add support for the OpenMP 5.0 task detach clause
2020-12-09 Kwok Cheung Yeung <k...@codesourcery.com>
gcc/
* builtin-types.def (BT_PTR_SIZED_INT): New primitive type.
(BT_FN_PSINT_VOID): New function type.
(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_PSINT):
...this. Add extra argument.
* gimplify.c (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_PSINT.
(BUILT_IN_GOMP_NEW_EVENT): New.
* omp-expand.c (expand_task_call): Add detach argument when generating
call to GOMP_task.
* omp-low.c (scan_sharing_clauses): Setup data environment for detach
clause.
(lower_detach_clause): New.
(lower_omp_taskreg): Call lower_detach_clause for detach clause. Add
Gimple statements generated for detach clause.
* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_DETACH.
* tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE_DETACH.
* tree.c (omp_clause_num_ops): Add entry for OMP_CLAUSE_DETACH.
(omp_clause_code_name): Add entry for OMP_CLAUSE_DETACH.
(walk_tree_1): Handle OMP_CLAUSE_DETACH.
* tree.h (OMP_CLAUSE_DETACH_EXPR): New.
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.
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.
* semantics.c (finish_omp_clauses): Handle OMP_CLAUSE_DETACH clause.
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.
* 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_PTR_SIZED_INT): New type.
(BT_FN_PSINT_VOID): New function type.
(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_PSINT):
...this. Add extra argument.
libgomp/
* fortran.c (omp_fulfill_event_): New.
* libgomp.h (struct gomp_allow_completion_event): New.
(struct gomp_task): Add detach_event field.
(struct gomp_team): Add task_detach_queue and task_detach_count
fields.
(gomp_finish_task): Delete detach_event.
* libgomp.map (OMP_5.0.1): Add omp_fulfill_event and omp_fulfill_event_.
(GOMP_5.0): Add GOMP_new_event.
* libgomp_g.h (GOMP_new_event): New.
(GOMP_task): Add uintptr_t 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_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_event field.
(GOMP_new_event): New.
(GOMP_task): Add detach argument. Initialize detach_event field.
Wait for detach event if task not deferred.
(task_fulfilled_p): New.
(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.
diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index 4a82ee4..d28ec75 100644
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
@@ -79,6 +79,7 @@ DEF_PRIMITIVE_TYPE (BT_UINT128, uint128_type_node
DEF_PRIMITIVE_TYPE (BT_WORD, (*lang_hooks.types.type_for_mode) (word_mode, 1))
DEF_PRIMITIVE_TYPE (BT_UNWINDWORD, (*lang_hooks.types.type_for_mode)
(targetm.unwind_word_mode (), 1))
+DEF_PRIMITIVE_TYPE (BT_PTR_SIZED_INT, pointer_sized_int_node)
DEF_PRIMITIVE_TYPE (BT_FLOAT, float_type_node)
DEF_PRIMITIVE_TYPE (BT_DOUBLE, double_type_node)
DEF_PRIMITIVE_TYPE (BT_LONGDOUBLE, long_double_type_node)
@@ -253,6 +254,7 @@ DEF_FUNCTION_TYPE_1 (BT_FN_LONG_LONGDOUBLE, BT_LONG,
BT_LONGDOUBLE)
DEF_FUNCTION_TYPE_1 (BT_FN_LONGLONG_FLOAT, BT_LONGLONG, BT_FLOAT)
DEF_FUNCTION_TYPE_1 (BT_FN_LONGLONG_DOUBLE, BT_LONGLONG, BT_DOUBLE)
DEF_FUNCTION_TYPE_1 (BT_FN_LONGLONG_LONGDOUBLE, BT_LONGLONG, BT_LONGDOUBLE)
+DEF_FUNCTION_TYPE_1 (BT_FN_PSINT_VOID, BT_PTR_SIZED_INT, BT_VOID)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTR, BT_VOID, BT_PTR)
DEF_FUNCTION_TYPE_1 (BT_FN_SIZE_CONST_STRING, BT_SIZE, BT_CONST_STRING)
DEF_FUNCTION_TYPE_1 (BT_FN_INT_CONST_STRING, BT_INT, BT_CONST_STRING)
@@ -756,10 +758,6 @@ DEF_FUNCTION_TYPE_8
(BT_FN_BOOL_UINT_ULLPTR_LONG_ULL_ULLPTR_ULLPTR_PTR_PTR,
BT_BOOL, BT_UINT, BT_PTR_ULONGLONG, BT_LONG, BT_ULONGLONG,
BT_PTR_ULONGLONG, BT_PTR_ULONGLONG, BT_PTR, BT_PTR)
-DEF_FUNCTION_TYPE_9 (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT,
- BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR,
- BT_PTR_FN_VOID_PTR_PTR, BT_LONG, BT_LONG,
- BT_BOOL, BT_UINT, BT_PTR, BT_INT)
DEF_FUNCTION_TYPE_9 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_PTR,
BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR,
BT_PTR, BT_PTR, BT_UINT, BT_PTR, BT_PTR)
@@ -767,6 +765,10 @@ DEF_FUNCTION_TYPE_9
(BT_FN_BOOL_LONG_LONG_LONG_LONG_LONG_LONGPTR_LONGPTR_PTR_PTR
BT_BOOL, BT_LONG, BT_LONG, BT_LONG, BT_LONG, BT_LONG,
BT_PTR_LONG, BT_PTR_LONG, BT_PTR, BT_PTR)
+DEF_FUNCTION_TYPE_10
(BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT_PSINT,
+ BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR,
+ BT_PTR_FN_VOID_PTR_PTR, BT_LONG, BT_LONG,
+ BT_BOOL, BT_UINT, BT_PTR, BT_INT, BT_PTR_SIZED_INT)
DEF_FUNCTION_TYPE_10
(BT_FN_BOOL_BOOL_ULL_ULL_ULL_LONG_ULL_ULLPTR_ULLPTR_PTR_PTR,
BT_BOOL, BT_BOOL, BT_ULONGLONG, BT_ULONGLONG,
BT_ULONGLONG, BT_LONG, BT_ULONGLONG, BT_PTR_ULONGLONG,
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index 5a493fe..fb784e9 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -94,6 +94,7 @@ enum pragma_omp_clause {
PRAGMA_OMP_CLAUSE_DEFAULT,
PRAGMA_OMP_CLAUSE_DEFAULTMAP,
PRAGMA_OMP_CLAUSE_DEPEND,
+ PRAGMA_OMP_CLAUSE_DETACH,
PRAGMA_OMP_CLAUSE_DEVICE,
PRAGMA_OMP_CLAUSE_DEVICE_TYPE,
PRAGMA_OMP_CLAUSE_DIST_SCHEDULE,
@@ -150,7 +151,6 @@ enum pragma_omp_clause {
PRAGMA_OACC_CLAUSE_COPYOUT,
PRAGMA_OACC_CLAUSE_CREATE,
PRAGMA_OACC_CLAUSE_DELETE,
- PRAGMA_OACC_CLAUSE_DETACH,
PRAGMA_OACC_CLAUSE_DEVICEPTR,
PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT,
PRAGMA_OACC_CLAUSE_FINALIZE,
@@ -173,6 +173,7 @@ enum pragma_omp_clause {
PRAGMA_OACC_CLAUSE_COPYIN = PRAGMA_OMP_CLAUSE_COPYIN,
PRAGMA_OACC_CLAUSE_DEVICE = PRAGMA_OMP_CLAUSE_DEVICE,
PRAGMA_OACC_CLAUSE_DEFAULT = PRAGMA_OMP_CLAUSE_DEFAULT,
+ PRAGMA_OACC_CLAUSE_DETACH = PRAGMA_OMP_CLAUSE_DETACH,
PRAGMA_OACC_CLAUSE_FIRSTPRIVATE = PRAGMA_OMP_CLAUSE_FIRSTPRIVATE,
PRAGMA_OACC_CLAUSE_IF = PRAGMA_OMP_CLAUSE_IF,
PRAGMA_OACC_CLAUSE_PRIVATE = PRAGMA_OMP_CLAUSE_PRIVATE,
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 7540a15..d3546d4 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15976,6 +15976,53 @@ c_parser_omp_clause_uniform (c_parser *parser, tree
list)
return list;
}
+/* OpenMP 5.0:
+ detach ( event-handle ) */
+
+static tree
+c_parser_omp_clause_detach (c_parser *parser, tree list)
+{
+ matching_parens parens;
+ location_t clause_loc = c_parser_peek_token (parser)->location;
+
+ if (!parens.require_open (parser))
+ return list;
+
+ if (c_parser_next_token_is_not (parser, CPP_NAME)
+ || c_parser_peek_token (parser)->id_kind != C_ID_ID)
+ {
+ c_parser_error (parser, "expected identifier");
+ return list;
+ }
+
+ tree t = lookup_name (c_parser_peek_token (parser)->value);
+ if (t == NULL_TREE)
+ {
+ undeclared_variable (c_parser_peek_token (parser)->location,
+ c_parser_peek_token (parser)->value);
+ return list;
+ }
+ c_parser_consume_token (parser);
+
+ tree type = TYPE_MAIN_VARIANT (TREE_TYPE (t));
+ if (!INTEGRAL_TYPE_P (type)
+ || TREE_CODE (type) != ENUMERAL_TYPE
+ || TYPE_NAME (type) != get_identifier ("omp_event_handle_t"))
+ {
+ error_at (clause_loc, "%<detach%> clause event handle "
+ "has type %qT rather than "
+ "%<omp_event_handle_t%>",
+ type);
+ return list;
+ }
+
+ tree u = build_omp_clause (clause_loc, OMP_CLAUSE_DETACH);
+ OMP_CLAUSE_DECL (u) = t;
+ OMP_CLAUSE_CHAIN (u) = list;
+ parens.skip_until_found_close (parser);
+ return u;
+}
+
/* Parse all OpenACC clauses. The set clauses allowed by the directive
is a bitmask in MASK. Return the list of clauses found. */
@@ -16242,6 +16289,10 @@ c_parser_omp_all_clauses (c_parser *parser,
omp_clause_mask mask,
clauses = c_parser_omp_clause_default (parser, clauses, false);
c_name = "default";
break;
+ case PRAGMA_OMP_CLAUSE_DETACH:
+ clauses = c_parser_omp_clause_detach (parser, clauses);
+ c_name = "detach";
+ break;
case PRAGMA_OMP_CLAUSE_FIRSTPRIVATE:
clauses = c_parser_omp_clause_firstprivate (parser, clauses);
c_name = "firstprivate";
@@ -19140,7 +19191,8 @@ c_parser_omp_single (location_t loc, c_parser *parser,
bool *if_p)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIORITY) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ALLOCATE) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION))
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DETACH))
static tree
c_parser_omp_task (location_t loc, c_parser *parser, bool *if_p)
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 286f3d9..0c2ed59 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -14942,6 +14942,11 @@ c_finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
pc = &OMP_CLAUSE_CHAIN (c);
continue;
+ case OMP_CLAUSE_DETACH:
+ t = OMP_CLAUSE_DECL (c);
+ pc = &OMP_CLAUSE_CHAIN (c);
+ continue;
+
case OMP_CLAUSE_IF:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_NUM_TEAMS:
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 8802124..5731a29 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -36957,6 +36957,66 @@ cp_parser_omp_clause_depend_sink (cp_parser *parser,
location_t clause_loc,
}
/* OpenMP 5.0:
+ detach ( event-handle ) */
+
+static tree
+cp_parser_omp_clause_detach (cp_parser *parser, tree list)
+{
+ matching_parens parens;
+
+ if (!parens.require_open (parser))
+ return list;
+
+ if (cp_lexer_next_token_is_not (parser->lexer, CPP_NAME))
+ {
+ cp_parser_error (parser, "expected identifier");
+ return list;
+ }
+
+ location_t id_loc = cp_lexer_peek_token (parser->lexer)->location;
+ tree t, identifier = cp_parser_identifier (parser);
+
+ if (identifier == error_mark_node)
+ t = error_mark_node;
+ else
+ {
+ t = cp_parser_lookup_name_simple
+ (parser, identifier,
+ cp_lexer_peek_token (parser->lexer)->location);
+ if (t == error_mark_node)
+ cp_parser_name_lookup_error (parser, identifier, t, NLE_NULL,
+ id_loc);
+ else
+ {
+ tree type = TYPE_MAIN_VARIANT (TREE_TYPE (t));
+ if (!INTEGRAL_TYPE_P (type)
+ || TREE_CODE (type) != ENUMERAL_TYPE
+ || DECL_NAME (TYPE_NAME (type))
+ != get_identifier ("omp_event_handle_t"))
+ {
+ error_at (id_loc, "%<detach%> clause event handle "
+ "has type %qT rather than "
+ "%<omp_event_handle_t%>",
+ type);
+ return list;
+ }
+ }
+ }
+
+ if (t == error_mark_node
+ || !parens.require_close (parser))
+ cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+
+ tree u = build_omp_clause (id_loc, OMP_CLAUSE_DETACH);
+ OMP_CLAUSE_DECL (u) = t;
+ OMP_CLAUSE_CHAIN (u) = list;
+
+ return u;
+}
+
+/* OpenMP 5.0:
iterators ( iterators-definition )
iterators-definition:
@@ -38012,6 +38072,10 @@ cp_parser_omp_all_clauses (cp_parser *parser,
omp_clause_mask mask,
token->location);
c_name = "depend";
break;
+ case PRAGMA_OMP_CLAUSE_DETACH:
+ clauses = cp_parser_omp_clause_detach (parser, clauses);
+ c_name = "detach";
+ break;
case PRAGMA_OMP_CLAUSE_MAP:
clauses = cp_parser_omp_clause_map (parser, clauses);
c_name = "map";
@@ -40546,7 +40610,8 @@ cp_parser_omp_single (cp_parser *parser, cp_token
*pragma_tok, bool *if_p)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIORITY) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ALLOCATE) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION))
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DETACH))
static tree
cp_parser_omp_task (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 5ff70ff..38eecf5 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -7394,6 +7394,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type
ort)
}
}
break;
+ case OMP_CLAUSE_DETACH:
+ t = OMP_CLAUSE_DECL (c);
+ break;
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_TO:
diff --git a/gcc/fortran/dump-parse-tree.c b/gcc/fortran/dump-parse-tree.c
index cab0fb2..baf1e3c 100644
--- a/gcc/fortran/dump-parse-tree.c
+++ b/gcc/fortran/dump-parse-tree.c
@@ -1700,6 +1700,12 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses)
show_expr (omp_clauses->priority);
fputc (')', dumpfile);
}
+ if (omp_clauses->detach)
+ {
+ fputs (" DETACH(", dumpfile);
+ show_expr (omp_clauses->detach);
+ fputc (')', dumpfile);
+ }
for (i = 0; i < OMP_IF_LAST; i++)
if (omp_clauses->if_exprs[i])
{
diff --git a/gcc/fortran/frontend-passes.c b/gcc/fortran/frontend-passes.c
index 83f6fd8..699b354 100644
--- a/gcc/fortran/frontend-passes.c
+++ b/gcc/fortran/frontend-passes.c
@@ -5597,6 +5597,7 @@ gfc_code_walker (gfc_code **c, walk_code_fn_t codefn,
walk_expr_fn_t exprfn,
WALK_SUBEXPR (co->ext.omp_clauses->hint);
WALK_SUBEXPR (co->ext.omp_clauses->num_tasks);
WALK_SUBEXPR (co->ext.omp_clauses->priority);
+ WALK_SUBEXPR (co->ext.omp_clauses->detach);
for (idx = 0; idx < OMP_IF_LAST; idx++)
WALK_SUBEXPR (co->ext.omp_clauses->if_exprs[idx]);
for (idx = 0;
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 6467985..fec96af 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1414,6 +1414,7 @@ typedef struct gfc_omp_clauses
struct gfc_expr *hint;
struct gfc_expr *num_tasks;
struct gfc_expr *priority;
+ struct gfc_expr *detach;
struct gfc_expr *if_exprs[OMP_IF_LAST];
enum gfc_omp_sched_kind dist_sched_kind;
struct gfc_expr *dist_chunk_size;
@@ -3102,6 +3103,7 @@ extern int gfc_default_character_kind;
extern int gfc_default_logical_kind;
extern int gfc_default_complex_kind;
extern int gfc_c_int_kind;
+extern int gfc_c_intptr_kind;
extern int gfc_atomic_int_kind;
extern int gfc_atomic_logical_kind;
extern int gfc_intio_kind;
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 68d0b65..4c58447 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -91,6 +91,7 @@ gfc_free_omp_clauses (gfc_omp_clauses *c)
gfc_free_expr (c->hint);
gfc_free_expr (c->num_tasks);
gfc_free_expr (c->priority);
+ gfc_free_expr (c->detach);
for (i = 0; i < OMP_IF_LAST; i++)
gfc_free_expr (c->if_exprs[i]);
gfc_free_expr (c->async_expr);
@@ -448,6 +449,39 @@ cleanup:
return MATCH_ERROR;
}
+/* Match detach(event-handle). */
+
+static match
+gfc_match_omp_detach (gfc_expr **expr)
+{
+ locus old_loc = gfc_current_locus;
+
+ if (gfc_match ("detach ( ") != MATCH_YES)
+ goto syntax_error;
+
+ if (gfc_match_variable (expr, 0) != MATCH_YES)
+ goto syntax_error;
+
+ if ((*expr)->ts.type != BT_INTEGER || (*expr)->ts.kind != gfc_c_intptr_kind)
+ {
+ gfc_error ("%qs at %L should be of type "
+ "integer(kind=omp_event_handle_kind)",
+ (*expr)->symtree->n.sym->name, &(*expr)->where);
+ return MATCH_ERROR;
+ }
+
+ if (gfc_match_char (')') != MATCH_YES)
+ goto syntax_error;
+
+ return MATCH_YES;
+
+syntax_error:
+ gfc_error ("Syntax error in OpenMP detach clause at %C");
+ gfc_current_locus = old_loc;
+ return MATCH_ERROR;
+
+}
+
/* Match depend(sink : ...) construct a namelist from it. */
static match
@@ -807,6 +841,7 @@ enum omp_mask1
OMP_CLAUSE_ATOMIC, /* OpenMP 5.0. */
OMP_CLAUSE_CAPTURE, /* OpenMP 5.0. */
OMP_CLAUSE_MEMORDER, /* OpenMP 5.0. */
+ OMP_CLAUSE_DETACH, /* OpenMP 5.0. */
OMP_CLAUSE_NOWAIT,
/* This must come last. */
OMP_MASK1_LAST
@@ -840,7 +875,6 @@ enum omp_mask2
OMP_CLAUSE_IF_PRESENT,
OMP_CLAUSE_FINALIZE,
OMP_CLAUSE_ATTACH,
- OMP_CLAUSE_DETACH,
/* This must come last. */
OMP_MASK2_LAST
};
@@ -1378,6 +1412,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const
omp_mask mask,
gfc_current_locus = old_loc;
}
if ((mask & OMP_CLAUSE_DETACH)
+ && !openacc
+ && gfc_match_omp_detach (&c->detach) == MATCH_YES)
+ continue;
+ if ((mask & OMP_CLAUSE_DETACH)
+ && openacc
&& gfc_match ("detach ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_DETACH, false,
@@ -2763,7 +2802,8 @@ cleanup:
(omp_mask (OMP_CLAUSE_PRIVATE) | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_SHARED | OMP_CLAUSE_IF | OMP_CLAUSE_DEFAULT \
| OMP_CLAUSE_UNTIED | OMP_CLAUSE_FINAL | OMP_CLAUSE_MERGEABLE \
- | OMP_CLAUSE_DEPEND | OMP_CLAUSE_PRIORITY | OMP_CLAUSE_IN_REDUCTION)
+ | OMP_CLAUSE_DEPEND | OMP_CLAUSE_PRIORITY | OMP_CLAUSE_IN_REDUCTION \
+ | OMP_CLAUSE_DETACH)
#define OMP_TASKLOOP_CLAUSES \
(omp_mask (OMP_CLAUSE_PRIVATE) | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_LASTPRIVATE | OMP_CLAUSE_SHARED | OMP_CLAUSE_IF
\
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 6b4ad6a..314e00d 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -3667,6 +3667,21 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
omp_clauses = gfc_trans_add_clause (c, omp_clauses);
}
+ if (clauses->detach)
+ {
+ tree detach;
+
+ gfc_init_se (&se, NULL);
+ gfc_conv_expr (&se, clauses->detach);
+ gfc_add_block_to_block (block, &se.pre);
+ detach = se.expr;
+ gfc_add_block_to_block (block, &se.post);
+
+ c = build_omp_clause (gfc_get_location (&where), OMP_CLAUSE_DETACH);
+ OMP_CLAUSE_DETACH_EXPR (c) = detach;
+ omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+ }
+
if (clauses->hint)
{
tree hint;
diff --git a/gcc/fortran/trans-types.c b/gcc/fortran/trans-types.c
index b7129dc..bd9ca14 100644
--- a/gcc/fortran/trans-types.c
+++ b/gcc/fortran/trans-types.c
@@ -114,6 +114,7 @@ int gfc_default_character_kind;
int gfc_default_logical_kind;
int gfc_default_complex_kind;
int gfc_c_int_kind;
+int gfc_c_intptr_kind;
int gfc_atomic_int_kind;
int gfc_atomic_logical_kind;
@@ -691,6 +692,8 @@ gfc_init_kinds (void)
/* Choose atomic kinds to match C's int. */
gfc_atomic_int_kind = gfc_c_int_kind;
gfc_atomic_logical_kind = gfc_c_int_kind;
+
+ gfc_c_intptr_kind = POINTER_SIZE / 8;
}
diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def
index 5736bba..7cac4dc 100644
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
@@ -53,6 +53,7 @@ DEF_PRIMITIVE_TYPE (BT_LONG, long_integer_type_node)
DEF_PRIMITIVE_TYPE (BT_ULONGLONG, long_long_unsigned_type_node)
DEF_PRIMITIVE_TYPE (BT_WORD, (*lang_hooks.types.type_for_mode) (word_mode, 1))
DEF_PRIMITIVE_TYPE (BT_SIZE, size_type_node)
+DEF_PRIMITIVE_TYPE (BT_PTR_SIZED_INT, pointer_sized_int_node)
DEF_PRIMITIVE_TYPE (BT_I1, builtin_type_for_size (BITS_PER_UNIT*1, 1))
DEF_PRIMITIVE_TYPE (BT_I2, builtin_type_for_size (BITS_PER_UNIT*2, 1))
@@ -86,6 +87,7 @@ DEF_FUNCTION_TYPE_1 (BT_FN_VOID_VPTR, BT_VOID,
BT_VOLATILE_PTR)
DEF_FUNCTION_TYPE_1 (BT_FN_INT_INT, BT_INT, BT_INT)
DEF_FUNCTION_TYPE_1 (BT_FN_UINT_UINT, BT_UINT, BT_UINT)
DEF_FUNCTION_TYPE_1 (BT_FN_PTR_PTR, BT_PTR, BT_PTR)
+DEF_FUNCTION_TYPE_1 (BT_FN_PSINT_VOID, BT_PTR_SIZED_INT, BT_VOID)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_INT, BT_VOID, BT_INT)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_BOOL, BT_VOID, BT_BOOL)
DEF_FUNCTION_TYPE_1 (BT_FN_BOOL_INT, BT_BOOL, BT_INT)
@@ -235,10 +237,6 @@ DEF_FUNCTION_TYPE_8
(BT_FN_BOOL_UINT_ULLPTR_LONG_ULL_ULLPTR_ULLPTR_PTR_PTR,
BT_BOOL, BT_UINT, BT_PTR_ULONGLONG, BT_LONG, BT_ULONGLONG,
BT_PTR_ULONGLONG, BT_PTR_ULONGLONG, BT_PTR, BT_PTR)
-DEF_FUNCTION_TYPE_9 (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT,
- BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR,
- BT_PTR_FN_VOID_PTR_PTR, BT_LONG, BT_LONG,
- BT_BOOL, BT_UINT, BT_PTR, BT_INT)
DEF_FUNCTION_TYPE_9 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_PTR,
BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR,
BT_PTR, BT_PTR, BT_UINT, BT_PTR, BT_PTR)
@@ -246,6 +244,10 @@ DEF_FUNCTION_TYPE_9
(BT_FN_BOOL_LONG_LONG_LONG_LONG_LONG_LONGPTR_LONGPTR_PTR_PTR
BT_BOOL, BT_LONG, BT_LONG, BT_LONG, BT_LONG, BT_LONG,
BT_PTR_LONG, BT_PTR_LONG, BT_PTR, BT_PTR)
+DEF_FUNCTION_TYPE_10
(BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT_PSINT,
+ BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR,
+ BT_PTR_FN_VOID_PTR_PTR, BT_LONG, BT_LONG,
+ BT_BOOL, BT_UINT, BT_PTR, BT_INT, BT_PTR_SIZED_INT)
DEF_FUNCTION_TYPE_10
(BT_FN_BOOL_BOOL_ULL_ULL_ULL_LONG_ULL_ULLPTR_ULLPTR_PTR_PTR,
BT_BOOL, BT_BOOL, BT_ULONGLONG, BT_ULONGLONG,
BT_ULONGLONG, BT_LONG, BT_ULONGLONG, BT_PTR_ULONGLONG,
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 53ec9ec..20d39b5 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -9733,6 +9733,19 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
}
break;
+ case OMP_CLAUSE_DETACH:
+ decl = OMP_CLAUSE_DECL (c);
+ if (outer_ctx)
+ {
+ splay_tree_node on
+ = splay_tree_lookup (outer_ctx->variables,
+ (splay_tree_key)decl);
+ if (on == NULL || (on->value & GOVD_DATA_SHARE_CLASS) == 0)
+ omp_firstprivatize_variable (outer_ctx, decl);
+ omp_notice_variable (outer_ctx, decl, true);
+ }
+ break;
+
case OMP_CLAUSE_IF:
if (OMP_CLAUSE_IF_MODIFIER (c) != ERROR_MARK
&& OMP_CLAUSE_IF_MODIFIER (c) != code)
@@ -10870,6 +10883,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p,
gimple_seq body, tree *list_p,
case OMP_CLAUSE_DEFAULTMAP:
case OMP_CLAUSE_ORDER:
case OMP_CLAUSE_BIND:
+ case OMP_CLAUSE_DETACH:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index f9b78ed..b3bf17f 100644
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -381,7 +381,7 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_PARALLEL_REDUCTIONS,
"GOMP_parallel_reductions",
BT_FN_UINT_OMPFN_PTR_UINT_UINT, ATTR_NOTHROW_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TASK, "GOMP_task",
- BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT,
+
BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT_PSINT,
ATTR_NOTHROW_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TASKLOOP, "GOMP_taskloop",
BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_UINT_LONG_INT_LONG_LONG_LONG,
@@ -451,3 +451,5 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_ALLOC,
ATTR_ALLOC_WARN_UNUSED_RESULT_SIZE_2_NOTHROW_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_FREE,
"GOMP_free", BT_FN_VOID_PTR_PTRMODE, ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_NEW_EVENT, "GOMP_new_event",
+ BT_FN_PSINT_VOID, ATTR_NOTHROW_LEAF_LIST)
diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
index c0e94e5..533f47d 100644
--- a/gcc/omp-expand.c
+++ b/gcc/omp-expand.c
@@ -762,6 +762,7 @@ expand_task_call (struct omp_region *region, basic_block bb,
tree depend = omp_find_clause (clauses, OMP_CLAUSE_DEPEND);
tree finalc = omp_find_clause (clauses, OMP_CLAUSE_FINAL);
tree priority = omp_find_clause (clauses, OMP_CLAUSE_PRIORITY);
+ tree detach = omp_find_clause (clauses, OMP_CLAUSE_DETACH);
unsigned int iflags
= (untied ? GOMP_TASK_FLAG_UNTIED : 0)
@@ -853,6 +854,11 @@ expand_task_call (struct omp_region *region, basic_block
bb,
priority = integer_zero_node;
gsi = gsi_last_nondebug_bb (bb);
+
+ detach = detach
+ ? fold_convert (pointer_sized_int_node, OMP_CLAUSE_DETACH_EXPR (detach))
+ : null_pointer_node;
+
tree t = gimple_omp_task_data_arg (entry_stmt);
if (t == NULL)
t2 = null_pointer_node;
@@ -875,10 +881,10 @@ expand_task_call (struct omp_region *region, basic_block
bb,
num_tasks, priority, startvar, endvar, step);
else
t = build_call_expr (builtin_decl_explicit (BUILT_IN_GOMP_TASK),
- 9, t1, t2, t3,
+ 10, t1, t2, t3,
gimple_omp_task_arg_size (entry_stmt),
gimple_omp_task_arg_align (entry_stmt), cond, flags,
- depend, priority);
+ depend, priority, detach);
force_gimple_operand_gsi (&gsi, t, true, NULL_TREE,
false, GSI_CONTINUE_LINKING);
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 09a8cbd..3dd15d1 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1412,6 +1412,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_NUM_GANGS:
case OMP_CLAUSE_NUM_WORKERS:
case OMP_CLAUSE_VECTOR_LENGTH:
+ case OMP_CLAUSE_DETACH:
if (ctx->outer)
scan_omp_op (&OMP_CLAUSE_OPERAND (c, 0), ctx->outer);
break;
@@ -1779,6 +1780,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_SIMDLEN:
case OMP_CLAUSE_ALIGNED:
case OMP_CLAUSE_DEPEND:
+ case OMP_CLAUSE_DETACH:
case OMP_CLAUSE_ALLOCATE:
case OMP_CLAUSE__LOOPTEMP_:
case OMP_CLAUSE__REDUCTEMP_:
@@ -11453,6 +11455,26 @@ create_task_copyfn (gomp_task *task_stmt, omp_context
*ctx)
}
static void
+lower_detach_clause (tree *pclauses, gimple_seq *iseq, omp_context *ctx)
+{
+ tree clause = omp_find_clause (*pclauses, OMP_CLAUSE_DETACH);
+ gcc_assert (clause);
+
+ tree event_decl = OMP_CLAUSE_DECL (clause);
+ tree event_ref = lookup_decl_in_outer_ctx (event_decl, ctx);
+ tree fn_decl = builtin_decl_explicit (BUILT_IN_GOMP_NEW_EVENT);
+ tree handle = create_tmp_var (pointer_sized_int_node);
+
+ gimple *call_stmt = gimple_build_call (fn_decl, 0);
+ gimple_call_set_lhs (call_stmt, handle);
+ gimple_seq_add_stmt (iseq, call_stmt);
+
+ gimplify_assign (event_ref,
+ fold_convert (TREE_TYPE (event_decl), handle),
+ iseq);
+}
+
+static void
lower_depend_clauses (tree *pclauses, gimple_seq *iseq, gimple_seq *oseq)
{
tree c, clauses;
@@ -11601,6 +11623,15 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
if (ws_num == 1)
gimple_omp_parallel_set_combined_p (stmt, true);
}
+
+ gimple_seq detach_ilist = NULL;
+ if (gimple_code (stmt) == GIMPLE_OMP_TASK
+ && omp_find_clause (clauses, OMP_CLAUSE_DETACH))
+ {
+ lower_detach_clause (gimple_omp_task_clauses_ptr (stmt), &detach_ilist,
+ ctx);
+ }
+
gimple_seq dep_ilist = NULL;
gimple_seq dep_olist = NULL;
if (gimple_code (stmt) == GIMPLE_OMP_TASK
@@ -11678,6 +11709,10 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
gimple_seq olist = NULL;
gimple_seq ilist = NULL;
+
+ if (detach_ilist)
+ gimple_seq_add_seq (&ilist, detach_ilist);
+
lower_send_clauses (clauses, &ilist, &olist, ctx);
lower_send_shared_vars (&ilist, &olist, ctx);
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 313a6af..5b028da 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -301,6 +301,9 @@ enum omp_clause_code {
/* OpenMP clause: to (variable-list). */
OMP_CLAUSE_TO,
+ /* OpenMP clause: detach (event-handle). */
+ OMP_CLAUSE_DETACH,
+
/* OpenACC clauses: {copy, copyin, copyout, create, delete, deviceptr,
device, host (self), present, present_or_copy (pcopy), present_or_copyin
(pcopyin), present_or_copyout (pcopyout), present_or_create (pcreate)}
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 5a93c4d..073d14b 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -1265,6 +1265,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int
spc, dump_flags_t flags)
case OMP_CLAUSE_FINALIZE:
pp_string (pp, "finalize");
break;
+ case OMP_CLAUSE_DETACH:
+ pp_string (pp, "detach(");
+ dump_generic_node (pp, OMP_CLAUSE_DETACH_EXPR (clause), spc, flags,
+ false);
+ pp_right_paren (pp);
+ break;
default:
gcc_unreachable ();
diff --git a/gcc/tree.c b/gcc/tree.c
index d6ba553..921c928 100644
--- a/gcc/tree.c
+++ b/gcc/tree.c
@@ -299,6 +299,7 @@ unsigned const char omp_clause_num_ops[] =
1, /* OMP_CLAUSE_LINK */
2, /* OMP_CLAUSE_FROM */
2, /* OMP_CLAUSE_TO */
+ 1, /* OMP_CLAUSE_DETACH */
2, /* OMP_CLAUSE_MAP */
1, /* OMP_CLAUSE_USE_DEVICE_PTR */
1, /* OMP_CLAUSE_USE_DEVICE_ADDR */
@@ -384,6 +385,7 @@ const char * const omp_clause_code_name[] =
"link",
"from",
"to",
+ "detach",
"map",
"use_device_ptr",
"use_device_addr",
@@ -12234,6 +12236,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE_HINT:
case OMP_CLAUSE_TO_DECLARE:
case OMP_CLAUSE_LINK:
+ case OMP_CLAUSE_DETACH:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
diff --git a/gcc/tree.h b/gcc/tree.h
index 078919b..c76ee74 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1588,6 +1588,9 @@ class auto_suppress_location_wrappers
#define OMP_CLAUSE_PRIORITY_EXPR(NODE) \
OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_PRIORITY),0)
+#define OMP_CLAUSE_DETACH_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DETACH),0)
+
/* OpenACC clause expressions */
#define OMP_CLAUSE_EXPR(NODE, CLAUSE) \
OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, CLAUSE), 0)
diff --git a/libgomp/fortran.c b/libgomp/fortran.c
index cd719f9..976b248 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 070d29c..88fc217 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -516,6 +516,12 @@ struct gomp_taskwait
gomp_sem_t taskwait_sem;
};
+struct gomp_allow_completion_event
+{
+ bool fulfilled;
+ gomp_sem_t completion_sem;
+};
+
/* This structure describes a "task" to be run by a thread. */
struct gomp_task
@@ -545,6 +551,8 @@ struct gomp_task
entries and the gomp_task in which they reside. */
struct priority_node pnode[3];
+ struct gomp_allow_completion_event *detach_event;
+
struct gomp_task_icv icv;
void (*fn) (void *);
void *fn_data;
@@ -685,6 +693,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[];
};
@@ -931,6 +943,8 @@ gomp_finish_task (struct gomp_task *task)
{
if (__builtin_expect (task->depend_hash != NULL, 0))
free (task->depend_hash);
+ if (task->detach_event)
+ free (task->detach_event);
}
/* team.c */
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index 2c95f78..434dfc3 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 {
@@ -347,6 +349,7 @@ GOMP_5.0 {
GOMP_loop_ull_nonmonotonic_runtime_start;
GOMP_loop_ull_ordered_start;
GOMP_loop_ull_start;
+ GOMP_new_event;
GOMP_parallel_loop_maybe_nonmonotonic_runtime;
GOMP_parallel_loop_nonmonotonic_runtime;
GOMP_parallel_reductions;
diff --git a/libgomp/libgomp_g.h b/libgomp/libgomp_g.h
index b20e186..ca27c53 100644
--- a/libgomp/libgomp_g.h
+++ b/libgomp/libgomp_g.h
@@ -293,8 +293,9 @@ extern bool GOMP_cancellation_point (int);
/* task.c */
+extern uintptr_t GOMP_new_event (void);
extern void GOMP_task (void (*) (void *), void *, void (*) (void *, void *),
- long, long, bool, unsigned, void **, int);
+ long, long, bool, unsigned, void **, int, uintptr_t);
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 4424a16..62b6c0f 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 3b7f0cb..7b70d8b 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,13 @@
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 eb1dcc4..5b4053f 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 9b8d2ff..0c6b556 100644
--- a/libgomp/priority_queue.c
+++ b/libgomp/priority_queue.c
@@ -168,6 +168,56 @@ priority_queue_verify (enum priority_queue_type type,
}
#endif /* _LIBGOMP_CHECKING_ */
+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;
+}
+
+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;
+}
+
+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 0ad78f5..c6fd80d 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 a95067c..ae1fcf7 100644
--- a/libgomp/task.c
+++ b/libgomp/task.c
@@ -29,6 +29,7 @@
#include "libgomp.h"
#include <stdlib.h>
#include <string.h>
+#include <stdio.h>
#include "gomp-constants.h"
typedef struct gomp_task_depend_entry *hash_entry_type;
@@ -86,6 +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_event = NULL;
}
/* Clean up a task, after completing it. */
@@ -326,6 +328,28 @@ gomp_task_handle_depend (struct gomp_task *task, struct
gomp_task *parent,
}
}
+uintptr_t
+GOMP_new_event ()
+{
+ struct gomp_allow_completion_event *event;
+
+ event = (struct gomp_allow_completion_event *)
+ gomp_malloc (sizeof (struct gomp_allow_completion_event));
+ event->fulfilled = false;
+ gomp_sem_init (&event->completion_sem, 0);
+
+ gomp_debug (0, "GOMP_new_event: %p\n", event);
+
+ return (uintptr_t) event;
+}
+
+static bool
+task_fulfilled_p (struct gomp_task *task)
+{
+ return __atomic_load_n (&task->detach_event->fulfilled,
+ __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,11 +371,14 @@ 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, uintptr_t detach)
{
struct gomp_thread *thr = gomp_thread ();
struct gomp_team *team = thr->ts.team;
+ struct gomp_allow_completion_event *detach_event =
+ detach ? (struct gomp_allow_completion_event *) detach : NULL;
+
#ifdef HAVE_BROKEN_POSIX_SEMAPHORES
/* If pthread_mutex_* is used for omp_*lock*, then each task must be
tied to one thread all the time. This means UNTIED tasks must be
@@ -404,6 +431,10 @@ 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_event = detach_event;
+
if (thr->task)
{
task.in_tied_task = thr->task->in_tied_task;
@@ -420,6 +451,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.detach_event->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)
@@ -435,6 +470,7 @@ GOMP_task (void (*fn) (void *), void *data, void (*cpyfn)
(void *, void *),
gomp_clear_parent (&task.children_queue);
gomp_mutex_unlock (&team->task_lock);
}
+
gomp_end_task ();
}
else
@@ -458,6 +494,8 @@ 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_event = detach_event;
thr->task = task;
if (cpyfn)
{
@@ -1325,6 +1363,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->detach_event);
+
+ 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 +1452,44 @@ 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_event
+ && !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->detach_event);
+ 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 +2401,22 @@ omp_in_final (void)
}
ialias (omp_in_final)
+
+void omp_fulfill_event(omp_event_handle_t event)
+{
+ struct gomp_allow_completion_event *ev =
+ (struct gomp_allow_completion_event *) event;
+ struct gomp_thread *thr = gomp_thread ();
+ struct gomp_team *team = thr ? thr->ts.team : NULL;
+
+ if (__atomic_load_n (&ev->fulfilled, __ATOMIC_RELAXED))
+ gomp_fatal ("omp_fulfill_event: Event already fulfilled!\n");
+
+ gomp_debug(0, "omp_fulfill_event: %p\n", ev);
+ __atomic_store_n (&ev->fulfilled, true, __ATOMIC_RELAXED);
+ gomp_sem_post (&ev->completion_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 cbc3aec..ee488f2 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..0be74af
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/task-detach-4.c
@@ -0,0 +1,22 @@
+/* { 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;
+
+ #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..ab80ed8
--- /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
+ {
+ #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..45683ee
--- /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
+ {
+ #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..96e6387
--- /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
+ !$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..54227ef
--- /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
+ !$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
