@@ -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)
@@ -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. */
@@ -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)
@@ -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)
@@ -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:
@@ -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:
@@ -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,
@@ -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
@@ -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 */
@@ -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)
@@ -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;
@@ -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 ();
@@ -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)
{
@@ -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;
+}