This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[gomp4.1] Depend clause support for offloading
- From: Jakub Jelinek <jakub at redhat dot com>
- To: Ilya Verbin <iverbin at gmail dot com>
- Cc: gcc-patches at gcc dot gnu dot org
- Date: Wed, 2 Sep 2015 17:58:54 +0200
- Subject: [gomp4.1] Depend clause support for offloading
- Authentication-results: sourceware.org; auth=none
- References: <20150731161610 dot GF1780 at tucnak dot redhat dot com> <20150828181335 dot GS9425 at tucnak dot redhat dot com> <20150831150753 dot GC1847 at tucnak dot redhat dot com> <20150902112114 dot GA19034 at msticlxl57 dot ims dot intel dot com>
- Reply-to: Jakub Jelinek <jakub at redhat dot com>
Hi!
On Wed, Sep 02, 2015 at 02:21:14PM +0300, Ilya Verbin wrote:
> On Mon, Aug 31, 2015 at 17:07:53 +0200, Jakub Jelinek wrote:
> > * gimplify.c (gimplify_scan_omp_clauses): Handle
> > struct element GOMP_MAP_FIRSTPRIVATE_POINTER.
>
> Have you seen this?
>
> gcc/gimplify.c: In function âvoid gimplify_scan_omp_clauses(tree_node**, gimple_statement_base**, omp_region_type, tree_code)â:
> gcc/gimplify.c:6578:12: error: âscâ may be used uninitialized in this function [-Werror=maybe-uninitialized]
> : *sc != c;
> ^
I haven't, but I haven't bootstrapped it for a while, just keep
doing make -C gcc -j16 -k check RUNTESTFLAGS=gomp.exp and
make check-target-libgomp. That said, this looks like a false positive,
but I've added a NULL initialization for it anyway.
Here is the start of the async offloading support I've talked about,
but nowait is not supported on the library side yet, only depend clause
(and for that I haven't added a testcase yet).
2015-09-02 Jakub Jelinek <jakub@redhat.com>
* gimplify.c (gimplify_scan_omp_clauses): Initialize sc
to NULL to avoid false positive warnings.
* omp-low.c (check_omp_nesting_restrictions): Diagnose
depend(source) or depend(sink:...) on #pragma omp target *.
(expand_omp_target): Pass flags and depend arguments to
GOMP_target_{41,update_41,enter_exit_data} libcalls.
(lower_depend_clauses): Change first argument from gimple
to tree * pointing to the stmt's clauses.
(lower_omp_taskreg): Adjust caller.
(lower_omp_target): Lower depend clauses. Always use 16-bit
kinds and 8 as align shift. Use
GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION for zero length array
section in map clause with delete kind.
* omp-builtins.def (BUILT_IN_GOMP_TARGET,
BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA): Add flags and depend arguments.
(BUILT_IN_GOMP_TARGET_UPDATE): Change library function name
to GOMP_target_update_41. Add flags and depend arguments,
remove unused argument.
* builtin-types.def (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR,
BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR): Remove.
(BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR,
BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR): New.
gcc/c/
* c-typeck.c (handle_omp_array_sections): Set
OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION even for
GOMP_MAP_DELETE kinds.
gcc/cp/
* semantics.c (handle_omp_array_sections): Set
OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION even for
GOMP_MAP_DELETE kinds.
gcc/fortran/
* types.def (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR,
BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR): Remove.
(BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR,
BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR): New.
include/
* gomp-constants.h (enum gomp_map_kind): Add
GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION.
(GOMP_TARGET_FLAG_NOWAIT, GOMP_TARGET_FLAG_EXIT_DATA): Define.
libgomp/
* libgomp_g.h (GOMP_target_41, GOMP_target_enter_exit_data): Add
flags and depend arguments.
(GOMP_target_update_41): New prototype.
* libgomp.h (gomp_task_maybe_wait_for_dependencies): New prototype.
* libgomp.map (GOMP_4.1): Add GOMP_target_update_41.
* task.c (gomp_task_maybe_wait_for_dependencies): Remove prototype.
No longer static.
* target.c (GOMP_target_41): Add flags and depend arguments. If
depend is non-NULL, wait until all dependencies are satisfied.
(GOMP_target_enter_exit_data): Likewise. Use
flags & GOMP_TARGET_FLAG_EXIT_DATA to determine if it is enter
or exit data construct, instead of analysing kinds.
(gomp_exit_data): Handle GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION.
(GOMP_target_update_41): New function.
* testsuite/libgomp.c/target-24.c: New test.
--- gcc/gimplify.c.jj 2015-08-31 16:57:23.000000000 +0200
+++ gcc/gimplify.c 2015-09-02 14:20:41.012253248 +0200
@@ -6557,8 +6557,8 @@ gimplify_scan_omp_clauses (tree *list_p,
}
else
{
- tree *osc = struct_map_to_clause->get (decl), *sc;
- tree *pt = NULL;
+ tree *osc = struct_map_to_clause->get (decl);
+ tree *sc = NULL, *pt = NULL;
if (!ptr && TREE_CODE (*osc) == TREE_LIST)
osc = &TREE_PURPOSE (*osc);
if (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS)
--- gcc/omp-low.c.jj 2015-09-01 17:39:05.000000000 +0200
+++ gcc/omp-low.c 2015-09-02 15:13:13.726567918 +0200
@@ -3440,6 +3440,19 @@ check_omp_nesting_restrictions (gimple s
}
break;
case GIMPLE_OMP_TARGET:
+ for (c = gimple_omp_target_clauses (stmt); c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
+ && (OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SOURCE
+ || OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SINK))
+ {
+ enum omp_clause_depend_kind kind = OMP_CLAUSE_DEPEND_KIND (c);
+ gcc_assert (kind == OMP_CLAUSE_DEPEND_SOURCE
+ || kind == OMP_CLAUSE_DEPEND_SINK);
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%<depend(%s)%> is only allowed in %<omp ordered%>",
+ kind == OMP_CLAUSE_DEPEND_SOURCE ? "source" : "sink");
+ return false;
+ }
for (; ctx != NULL; ctx = ctx->outer)
{
if (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET)
@@ -10639,9 +10652,10 @@ expand_omp_target (struct omp_region *re
/* Emit a library call to launch the offloading region, or do data
transfers. */
- tree t1, t2, t3, t4, device, cond, c, clauses;
+ tree t1, t2, t3, t4, device, cond, depend, c, clauses;
enum built_in_function start_ix;
location_t clause_loc;
+ unsigned int flags_i = 0;
switch (gimple_omp_target_kind (entry_stmt))
{
@@ -10655,8 +10669,11 @@ expand_omp_target (struct omp_region *re
start_ix = BUILT_IN_GOMP_TARGET_UPDATE;
break;
case GF_OMP_TARGET_KIND_ENTER_DATA:
+ start_ix = BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA;
+ break;
case GF_OMP_TARGET_KIND_EXIT_DATA:
start_ix = BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA;
+ flags_i |= GOMP_TARGET_FLAG_EXIT_DATA;
break;
case GF_OMP_TARGET_KIND_OACC_PARALLEL:
case GF_OMP_TARGET_KIND_OACC_KERNELS:
@@ -10702,6 +10719,10 @@ expand_omp_target (struct omp_region *re
else
clause_loc = gimple_location (entry_stmt);
+ c = find_omp_clause (clauses, OMP_CLAUSE_NOWAIT);
+ if (c)
+ flags_i |= GOMP_TARGET_FLAG_NOWAIT;
+
/* Ensure 'device' is of the correct type. */
device = fold_convert_loc (clause_loc, integer_type_node, device);
@@ -10781,10 +10802,6 @@ expand_omp_target (struct omp_region *re
args.quick_push (device);
if (offloaded)
args.quick_push (build_fold_addr_expr (child_fn));
- /* This const void * is part of the current ABI, but we're not actually using
- it. */
- if (start_ix == BUILT_IN_GOMP_TARGET_UPDATE)
- args.quick_push (build_zero_cst (ptr_type_node));
args.quick_push (t1);
args.quick_push (t2);
args.quick_push (t3);
@@ -10792,10 +10809,18 @@ expand_omp_target (struct omp_region *re
switch (start_ix)
{
case BUILT_IN_GOACC_DATA_START:
- case BUILT_IN_GOMP_TARGET:
case BUILT_IN_GOMP_TARGET_DATA:
+ break;
+ case BUILT_IN_GOMP_TARGET:
case BUILT_IN_GOMP_TARGET_UPDATE:
case BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA:
+ args.quick_push (build_int_cst (unsigned_type_node, flags_i));
+ c = find_omp_clause (clauses, OMP_CLAUSE_DEPEND);
+ if (c)
+ depend = OMP_CLAUSE_DECL (c);
+ else
+ depend = build_int_cst (ptr_type_node, 0);
+ args.quick_push (depend);
break;
case BUILT_IN_GOACC_PARALLEL:
{
@@ -10891,8 +10916,7 @@ expand_omp_target (struct omp_region *re
gcc_assert (g && gimple_code (g) == GIMPLE_OMP_TARGET);
gsi_remove (&gsi, true);
}
- if (data_region
- && region->exit)
+ if (data_region && region->exit)
{
gsi = gsi_last_bb (region->exit);
g = gsi_stmt (gsi);
@@ -12923,14 +12947,13 @@ create_task_copyfn (gomp_task *task_stmt
}
static void
-lower_depend_clauses (gimple stmt, gimple_seq *iseq, gimple_seq *oseq)
+lower_depend_clauses (tree *pclauses, gimple_seq *iseq, gimple_seq *oseq)
{
tree c, clauses;
gimple g;
size_t n_in = 0, n_out = 0, idx = 2, i;
- clauses = find_omp_clause (gimple_omp_task_clauses (stmt),
- OMP_CLAUSE_DEPEND);
+ clauses = find_omp_clause (*pclauses, OMP_CLAUSE_DEPEND);
gcc_assert (clauses);
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND)
@@ -12977,11 +13000,10 @@ lower_depend_clauses (gimple stmt, gimpl
gimple_seq_add_stmt (iseq, g);
}
}
- tree *p = gimple_omp_task_clauses_ptr (stmt);
c = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE_DEPEND);
OMP_CLAUSE_DECL (c) = build_fold_addr_expr (array);
- OMP_CLAUSE_CHAIN (c) = *p;
- *p = c;
+ OMP_CLAUSE_CHAIN (c) = *pclauses;
+ *pclauses = c;
tree clobber = build_constructor (type, NULL);
TREE_THIS_VOLATILE (clobber) = 1;
g = gimple_build_assign (array, clobber);
@@ -13026,7 +13048,8 @@ lower_omp_taskreg (gimple_stmt_iterator
{
push_gimplify_context ();
dep_bind = gimple_build_bind (NULL, NULL, make_node (BLOCK));
- lower_depend_clauses (stmt, &dep_ilist, &dep_olist);
+ lower_depend_clauses (gimple_omp_task_clauses_ptr (stmt),
+ &dep_ilist, &dep_olist);
}
if (ctx->srecord_type)
@@ -13124,7 +13147,7 @@ lower_omp_target (gimple_stmt_iterator *
tree clauses;
tree child_fn, t, c;
gomp_target *stmt = as_a <gomp_target *> (gsi_stmt (*gsi_p));
- gbind *tgt_bind, *bind;
+ gbind *tgt_bind, *bind, *dep_bind = NULL;
gimple_seq tgt_body, olist, ilist, orlist, irlist, new_body;
location_t loc = gimple_location (stmt);
bool offloaded, data_region;
@@ -13153,6 +13176,16 @@ lower_omp_target (gimple_stmt_iterator *
clauses = gimple_omp_target_clauses (stmt);
+ gimple_seq dep_ilist = NULL;
+ gimple_seq dep_olist = NULL;
+ if (find_omp_clause (clauses, OMP_CLAUSE_DEPEND))
+ {
+ push_gimplify_context ();
+ dep_bind = gimple_build_bind (NULL, NULL, make_node (BLOCK));
+ lower_depend_clauses (gimple_omp_task_clauses_ptr (stmt),
+ &dep_ilist, &dep_olist);
+ }
+
tgt_bind = NULL;
tgt_body = NULL;
if (offloaded)
@@ -13378,19 +13411,8 @@ lower_omp_target (gimple_stmt_iterator *
DECL_NAMELESS (TREE_VEC_ELT (t, 1)) = 1;
TREE_ADDRESSABLE (TREE_VEC_ELT (t, 1)) = 1;
TREE_STATIC (TREE_VEC_ELT (t, 1)) = 1;
- tree tkind_type;
- int talign_shift;
- if (is_gimple_omp_oacc (stmt)
- || gimple_omp_target_kind (stmt) != GF_OMP_TARGET_KIND_UPDATE)
- {
- tkind_type = short_unsigned_type_node;
- talign_shift = 8;
- }
- else
- {
- tkind_type = unsigned_char_type_node;
- talign_shift = 3;
- }
+ tree tkind_type = short_unsigned_type_node;
+ int talign_shift = 8;
TREE_VEC_ELT (t, 2)
= create_tmp_var (build_array_type_nelts (tkind_type, map_cnt),
".omp_data_kinds");
@@ -13550,6 +13572,8 @@ lower_omp_target (gimple_stmt_iterator *
case GOMP_MAP_RELEASE:
tkind_zero = GOMP_MAP_ZERO_LEN_ARRAY_SECTION;
break;
+ case GOMP_MAP_DELETE:
+ tkind_zero = GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION;
default:
break;
}
@@ -14039,7 +14063,7 @@ lower_omp_target (gimple_stmt_iterator *
bind = gimple_build_bind (NULL, NULL,
tgt_bind ? gimple_bind_block (tgt_bind)
: NULL_TREE);
- gsi_replace (gsi_p, bind, true);
+ gsi_replace (gsi_p, dep_bind ? dep_bind : bind, true);
gimple_bind_add_seq (bind, irlist);
gimple_bind_add_seq (bind, ilist);
gimple_bind_add_stmt (bind, stmt);
@@ -14047,6 +14071,14 @@ lower_omp_target (gimple_stmt_iterator *
gimple_bind_add_seq (bind, orlist);
pop_gimplify_context (NULL);
+
+ if (dep_bind)
+ {
+ gimple_bind_add_seq (dep_bind, dep_ilist);
+ gimple_bind_add_stmt (dep_bind, bind);
+ gimple_bind_add_seq (dep_bind, dep_olist);
+ pop_gimplify_context (dep_bind);
+ }
}
/* Expand code for an OpenMP teams directive. */
--- gcc/omp-builtins.def.jj 2015-06-18 15:24:31.000000000 +0200
+++ gcc/omp-builtins.def 2015-09-02 12:51:00.710561827 +0200
@@ -263,15 +263,17 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_SINGLE_C
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_SINGLE_COPY_END, "GOMP_single_copy_end",
BT_FN_VOID_PTR, ATTR_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET, "GOMP_target_41",
- BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+ BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR,
+ ATTR_NOTHROW_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_DATA, "GOMP_target_data_41",
BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_END_DATA, "GOMP_target_end_data",
BT_FN_VOID, ATTR_NOTHROW_LIST)
-DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_UPDATE, "GOMP_target_update",
- BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_UPDATE, "GOMP_target_update_41",
+ BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR,
+ ATTR_NOTHROW_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA,
"GOMP_target_enter_exit_data",
- BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+ BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR, ATTR_NOTHROW_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TEAMS, "GOMP_teams",
BT_FN_VOID_UINT_UINT, ATTR_NOTHROW_LIST)
--- gcc/builtin-types.def.jj 2015-06-18 15:24:31.000000000 +0200
+++ gcc/builtin-types.def 2015-09-02 12:51:51.201829660 +0200
@@ -524,11 +524,6 @@ DEF_FUNCTION_TYPE_6 (BT_FN_BOOL_VPTR_PTR
BT_INT)
DEF_FUNCTION_TYPE_6 (BT_FN_BOOL_SIZE_VPTR_PTR_PTR_INT_INT, BT_BOOL, BT_SIZE,
BT_VOLATILE_PTR, BT_PTR, BT_PTR, BT_INT, BT_INT)
-DEF_FUNCTION_TYPE_6 (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR,
- BT_VOID, BT_INT, BT_PTR, BT_SIZE, BT_PTR, BT_PTR, BT_PTR)
-DEF_FUNCTION_TYPE_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR,
- BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR,
- BT_PTR, BT_PTR)
DEF_FUNCTION_TYPE_7 (BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_UINT,
BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT,
@@ -537,7 +532,13 @@ DEF_FUNCTION_TYPE_7 (BT_FN_BOOL_BOOL_ULL
BT_BOOL, BT_BOOL, BT_ULONGLONG, BT_ULONGLONG,
BT_ULONGLONG, BT_ULONGLONG,
BT_PTR_ULONGLONG, BT_PTR_ULONGLONG)
+DEF_FUNCTION_TYPE_7 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR,
+ BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR, BT_UINT,
+ BT_PTR)
+DEF_FUNCTION_TYPE_8 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR,
+ BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR,
+ BT_PTR, BT_PTR, BT_UINT, BT_PTR)
DEF_FUNCTION_TYPE_8 (BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_LONG_UINT,
BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT,
BT_LONG, BT_LONG, BT_LONG, BT_LONG, BT_UINT)
--- gcc/c/c-typeck.c.jj 2015-08-31 16:57:23.000000000 +0200
+++ gcc/c/c-typeck.c 2015-09-02 13:53:39.487580457 +0200
@@ -12070,6 +12070,7 @@ handle_omp_array_sections (tree c, bool
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_RELEASE:
+ case GOMP_MAP_DELETE:
OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
break;
default:
--- gcc/cp/semantics.c.jj 2015-08-31 16:57:23.000000000 +0200
+++ gcc/cp/semantics.c 2015-09-02 13:54:11.019128248 +0200
@@ -4869,6 +4869,7 @@ handle_omp_array_sections (tree c, bool
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_RELEASE:
+ case GOMP_MAP_DELETE:
OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
break;
default:
--- gcc/fortran/types.def.jj 2015-06-18 15:24:31.000000000 +0200
+++ gcc/fortran/types.def 2015-09-02 12:52:20.089410765 +0200
@@ -189,11 +189,6 @@ DEF_FUNCTION_TYPE_6 (BT_FN_BOOL_VPTR_PTR
BT_INT)
DEF_FUNCTION_TYPE_6 (BT_FN_BOOL_SIZE_VPTR_PTR_PTR_INT_INT, BT_BOOL, BT_SIZE,
BT_VOLATILE_PTR, BT_PTR, BT_PTR, BT_INT, BT_INT)
-DEF_FUNCTION_TYPE_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR,
- BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR,
- BT_PTR, BT_PTR)
-DEF_FUNCTION_TYPE_6 (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR,
- BT_VOID, BT_INT, BT_PTR, BT_SIZE, BT_PTR, BT_PTR, BT_PTR)
DEF_FUNCTION_TYPE_7 (BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_UINT,
BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT,
@@ -202,10 +197,16 @@ DEF_FUNCTION_TYPE_7 (BT_FN_BOOL_BOOL_ULL
BT_BOOL, BT_BOOL, BT_ULONGLONG, BT_ULONGLONG,
BT_ULONGLONG, BT_ULONGLONG,
BT_PTR_ULONGLONG, BT_PTR_ULONGLONG)
+DEF_FUNCTION_TYPE_7 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR,
+ BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR, BT_UINT,
+ BT_PTR)
DEF_FUNCTION_TYPE_8 (BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_LONG_UINT,
BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT,
BT_LONG, BT_LONG, BT_LONG, BT_LONG, BT_UINT)
+DEF_FUNCTION_TYPE_8 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR,
+ BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR,
+ BT_PTR, BT_PTR, BT_UINT, 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,
--- include/gomp-constants.h.jj 2015-07-31 16:55:38.000000000 +0200
+++ include/gomp-constants.h 2015-09-02 13:53:09.065016663 +0200
@@ -110,6 +110,10 @@ enum gomp_map_kind
(address of the last adjacent entry plus its size). */
GOMP_MAP_STRUCT = (GOMP_MAP_FLAG_ALWAYS
| GOMP_MAP_FLAG_SPECIAL | 0),
+ /* Forced deallocation of zero length array section. */
+ GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
+ = (GOMP_MAP_FLAG_ALWAYS
+ | GOMP_MAP_FLAG_SPECIAL | 3),
/* OpenMP 4.1 alias for forced deallocation. */
GOMP_MAP_DELETE = GOMP_MAP_FORCE_DEALLOC,
/* Decrement usage count and deallocate if zero. */
@@ -171,4 +175,8 @@ enum gomp_map_kind
#define GOMP_TASK_FLAG_IF (1 << 10)
#define GOMP_TASK_FLAG_NOGROUP (1 << 11)
+/* GOMP_target{_41,update_41,enter_exit_data} flags argument. */
+#define GOMP_TARGET_FLAG_NOWAIT (1 << 0)
+#define GOMP_TARGET_FLAG_EXIT_DATA (1 << 1)
+
#endif
--- libgomp/libgomp_g.h.jj 2015-06-18 15:24:32.000000000 +0200
+++ libgomp/libgomp_g.h 2015-09-02 12:50:21.794126150 +0200
@@ -217,7 +217,7 @@ extern void GOMP_single_copy_end (void *
extern void GOMP_target (int, void (*) (void *), const void *,
size_t, void **, size_t *, unsigned char *);
extern void GOMP_target_41 (int, void (*) (void *), size_t, void **, size_t *,
- unsigned short *);
+ unsigned short *, unsigned int, void **);
extern void GOMP_target_data (int, const void *,
size_t, void **, size_t *, unsigned char *);
extern void GOMP_target_data_41 (int, size_t, void **, size_t *,
@@ -225,8 +225,11 @@ extern void GOMP_target_data_41 (int, si
extern void GOMP_target_end_data (void);
extern void GOMP_target_update (int, const void *,
size_t, void **, size_t *, unsigned char *);
+extern void GOMP_target_update_41 (int, size_t, void **, size_t *,
+ unsigned short *, unsigned int, void **);
extern void GOMP_target_enter_exit_data (int, size_t, void **, size_t *,
- unsigned short *);
+ unsigned short *, unsigned int,
+ void **);
extern void GOMP_teams (unsigned int, unsigned int);
/* oacc-parallel.c */
--- libgomp/libgomp.h.jj 2015-08-31 16:54:12.000000000 +0200
+++ libgomp/libgomp.h 2015-09-02 15:21:44.722166933 +0200
@@ -650,6 +650,7 @@ extern void gomp_init_task (struct gomp_
struct gomp_task_icv *);
extern void gomp_end_task (void);
extern void gomp_barrier_handle_tasks (gomp_barrier_state_t);
+extern void gomp_task_maybe_wait_for_dependencies (void **);
static void inline
gomp_finish_task (struct gomp_task *task)
--- libgomp/libgomp.map.jj 2015-07-10 18:49:17.000000000 +0200
+++ libgomp/libgomp.map 2015-09-02 12:01:18.132047752 +0200
@@ -268,6 +268,7 @@ GOMP_4.1 {
global:
GOMP_target_41;
GOMP_target_data_41;
+ GOMP_target_update_41;
GOMP_target_enter_exit_data;
GOMP_taskloop;
GOMP_taskloop_ull;
--- libgomp/task.c.jj 2015-08-31 16:54:12.000000000 +0200
+++ libgomp/task.c 2015-09-02 15:22:14.162740580 +0200
@@ -108,8 +108,6 @@ gomp_clear_parent (struct gomp_task *chi
while (task != children);
}
-static void gomp_task_maybe_wait_for_dependencies (void **depend);
-
/* 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.
@@ -987,7 +985,7 @@ GOMP_taskwait (void)
DEPEND is as in GOMP_task. */
-static void
+void
gomp_task_maybe_wait_for_dependencies (void **depend)
{
struct gomp_thread *thr = gomp_thread ();
--- libgomp/target.c.jj 2015-08-31 16:57:23.000000000 +0200
+++ libgomp/target.c 2015-09-02 15:30:23.350656259 +0200
@@ -1247,10 +1247,22 @@ GOMP_target (int device, void (*fn) (voi
void
GOMP_target_41 (int device, void (*fn) (void *), size_t mapnum,
- void **hostaddrs, size_t *sizes, unsigned short *kinds)
+ void **hostaddrs, size_t *sizes, unsigned short *kinds,
+ unsigned int flags, void **depend)
{
struct gomp_device_descr *devicep = resolve_device (device);
+ /* If there are depend clauses, but nowait is not present,
+ block the parent task until the dependencies are resolved
+ and then just continue with the rest of the function as if it
+ is a merged task. */
+ if (depend != NULL)
+ {
+ struct gomp_thread *thr = gomp_thread ();
+ if (thr->task && thr->task->depend_hash)
+ gomp_task_maybe_wait_for_dependencies (depend);
+ }
+
if (devicep == NULL
|| !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
{
@@ -1386,6 +1398,31 @@ GOMP_target_update (int device, const vo
gomp_update (devicep, mapnum, hostaddrs, sizes, kinds, false);
}
+void
+GOMP_target_update_41 (int device, size_t mapnum, void **hostaddrs,
+ size_t *sizes, unsigned short *kinds,
+ unsigned int flags, void **depend)
+{
+ struct gomp_device_descr *devicep = resolve_device (device);
+
+ /* If there are depend clauses, but nowait is not present,
+ block the parent task until the dependencies are resolved
+ and then just continue with the rest of the function as if it
+ is a merged task. */
+ if (depend != NULL)
+ {
+ struct gomp_thread *thr = gomp_thread ();
+ if (thr->task && thr->task->depend_hash)
+ gomp_task_maybe_wait_for_dependencies (depend);
+ }
+
+ if (devicep == NULL
+ || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
+ return;
+
+ gomp_update (devicep, mapnum, hostaddrs, sizes, kinds, true);
+}
+
static void
gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum,
void **hostaddrs, size_t *sizes, unsigned short *kinds)
@@ -1404,9 +1441,11 @@ gomp_exit_data (struct gomp_device_descr
case GOMP_MAP_DELETE:
case GOMP_MAP_RELEASE:
case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
+ case GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION:
cur_node.host_start = (uintptr_t) hostaddrs[i];
cur_node.host_end = cur_node.host_start + sizes[i];
- splay_tree_key k = kind == GOMP_MAP_ZERO_LEN_ARRAY_SECTION
+ splay_tree_key k = (kind == GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
+ || kind == GOMP_MAP_ZERO_LEN_ARRAY_SECTION)
? gomp_map_lookup (&devicep->mem_map, &cur_node)
: splay_tree_lookup (&devicep->mem_map, &cur_node);
if (!k)
@@ -1414,7 +1453,9 @@ gomp_exit_data (struct gomp_device_descr
if (k->refcount > 0 && k->refcount != REFCOUNT_INFINITY)
k->refcount--;
- if (kind == GOMP_MAP_DELETE && k->refcount != REFCOUNT_INFINITY)
+ if ((kind == GOMP_MAP_DELETE
+ || kind == GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION)
+ && k->refcount != REFCOUNT_INFINITY)
k->refcount = 0;
if ((kind == GOMP_MAP_FROM && k->refcount == 0)
@@ -1447,42 +1488,28 @@ gomp_exit_data (struct gomp_device_descr
void
GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
- size_t *sizes, unsigned short *kinds)
+ size_t *sizes, unsigned short *kinds,
+ unsigned int flags, void **depend)
{
struct gomp_device_descr *devicep = resolve_device (device);
+ /* If there are depend clauses, but nowait is not present,
+ block the parent task until the dependencies are resolved
+ and then just continue with the rest of the function as if it
+ is a merged task. */
+ if (depend != NULL)
+ {
+ struct gomp_thread *thr = gomp_thread ();
+ if (thr->task && thr->task->depend_hash)
+ gomp_task_maybe_wait_for_dependencies (depend);
+ }
+
if (devicep == NULL
|| !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
return;
- /* Determine if this is an "omp target enter data". */
- const int typemask = 0xff;
- bool is_enter_data = false;
size_t i;
- for (i = 0; i < mapnum; i++)
- {
- unsigned char kind = kinds[i] & typemask;
-
- if (kind == GOMP_MAP_ALLOC
- || kind == GOMP_MAP_TO
- || kind == GOMP_MAP_ALWAYS_TO
- || kind == GOMP_MAP_STRUCT)
- {
- is_enter_data = true;
- break;
- }
-
- if (kind == GOMP_MAP_FROM
- || kind == GOMP_MAP_ALWAYS_FROM
- || kind == GOMP_MAP_DELETE
- || kind == GOMP_MAP_RELEASE
- || kind == GOMP_MAP_ZERO_LEN_ARRAY_SECTION)
- break;
-
- gomp_fatal ("GOMP_target_enter_exit_data unhandled kind 0x%.2x", kind);
- }
-
- if (is_enter_data)
+ if ((flags & GOMP_TARGET_FLAG_EXIT_DATA) == 0)
for (i = 0; i < mapnum; i++)
if ((kinds[i] & 0xff) == GOMP_MAP_STRUCT)
{
--- libgomp/testsuite/libgomp.c/target-24.c.jj 2015-09-02 16:52:08.540815330 +0200
+++ libgomp/testsuite/libgomp.c/target-24.c 2015-09-02 16:54:13.176019999 +0200
@@ -0,0 +1,43 @@
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+ int d = omp_get_default_device ();
+ int id = omp_get_initial_device ();
+
+ if (d < 0 || d >= omp_get_num_devices ())
+ d = id;
+
+ int a[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 };
+ int *b = a;
+ int shared_mem = 0;
+ #pragma omp target map (alloc: shared_mem)
+ shared_mem = 1;
+ if (omp_target_is_present (b, 0, d) != shared_mem)
+ abort ();
+ #pragma omp target enter data map (to: a)
+ if (omp_target_is_present (b, 0, d) == 0)
+ abort ();
+ #pragma omp target enter data map (alloc: b[:0])
+ if (omp_target_is_present (b, 0, d) == 0)
+ abort ();
+ #pragma omp target exit data map (release: b[:0])
+ if (omp_target_is_present (b, 0, d) == 0)
+ abort ();
+ #pragma omp target exit data map (release: b[:0])
+ if (omp_target_is_present (b, 0, d) != shared_mem)
+ abort ();
+ #pragma omp target enter data map (to: a)
+ if (omp_target_is_present (b, 0, d) == 0)
+ abort ();
+ #pragma omp target enter data map (always, to: b[:0])
+ if (omp_target_is_present (b, 0, d) == 0)
+ abort ();
+ #pragma omp target exit data map (delete: b[:0])
+ if (omp_target_is_present (b, 0, d) != shared_mem)
+ abort ();
+ #pragma omp target exit data map (from: b[:0])
+ return 0;
+}
Jakub