Hi Jakub,
this patch set implements more of OpenMP 5.0 mapping, specifically this part in
2.19.7.1 map Clause:
"If a list item in a map clause is a structure element then all other structure
elements of the
containing structure variable form a structure sibling list. The map clause
and the structure sibling
list are associated with the same construct. If a corresponding list item of
the structure sibling list
item is present in the device data environment when the construct is
encountered then:
* If the 1 structure sibling list item does not appear in a map clause on the
construct then:
– If the construct is a target, target data, or target enter data construct
then the
structure sibling list item is treated as if it is a list item in a map
clause on the construct with a
map-type of alloc.
– If the construct is target exit data construct, then the structure sibling
list item is treated
as if it is a list item in a map clause on the construct with a map-type of
release."
While really wordy, I believe this simply means that maps of structure element
fields have their reference
counts increased/decreased in a uniform fashion, i.e. the are alloc/release'd
together, instead of
having parts of the structure possibly deallocated while others still exist on
the device.
In general, upon encountering a construct, we can't statically determine and
insert alloc/release maps
for each element of a structure variable, since we don't really know which
region of the structure is
currently mapped or not, hence this probably can't be properly implemented in
the compiler.
Instead this patch tries to do the equivalent in the runtime: I've modified the
handling of the
(GOMP_MAP_STRUCT, <field-map1>, <field-map2>, ...) sequence to:
(1) Create just a single splay_tree_key to represent the entire structure's
mapped-region
(all element target_var_desc's now reference this same key instead of
creating their own), and
(2) Associated an increment/decrement of the splay_tree_key refcount only
with the leading GOMP_MAP_STRUCT,
not with each individual struct element, e.g. those element
tgt_var_desc's mainly are now only used to
execute the host<->device copying logic.
This implies that, GOMP_MAP_STRUCT is needed also in "exit data" directives
too, so a small patch in
gcc/gimplify.c has been made to NOT remove this map for OpenMP target exit
data. OpenACC has not been touched.
(There are some parts of the libgomp changes with are related to
differentiating OpenMP/OpenACC cases,
which also exists in the last 5.0-mapping patch, also included here for
self-completeness)
This patch contains three libgomp testcases, the first one of which was also
included in the last 5.0-mapping
patch set as an XFAIL, but now passes with this patch.
Tobias' had an earlier issue with Fortran arrays (I forgot which kind it was
called) where sub-struct fields
clashed with each other, causing a libgomp runtime mapping fail. That problem
should be fixed with this
patch, since the capturing of all tgt_var_desc's key references into a single
splay_tree_key inherently avoids
the multiple overlapping key behavior. I have lightly tested Tobias' testcase
he gave me earlier on this,
and this part of the issue appears to be solved, however it still needs the
first 5.0-mapping patch combined
with this patch to completely work, since the Fortran array struct needs
pointer-attachment/detachment of the
data to really work before and after the target region.
This patch has been tested on x86_64-linux with nvptx offloading with no
regressions, also currently testing
for powerpc64le-linux, seeking approval for trunk.
(BTW Jakub, thanks for your review of the other first patch set, I will be
working on that revision next).
Thanks,
Chung-Lin
2020-10-23 Chung-Lin Tang <clt...@codesourcery.com>
gcc/
* gimplify.c (gimplify_adjust_omp_clauses): Do not remove
GOMP_MAP_STRUCT clauses for OpenMP target exit data constructs.
libgomp/
* libgomp.h (struct target_var_desc): New 'bool is_struct' field,
update comments for 'length' field.
(enum gomp_map_vars_kind): Adjust enum values to be bit-flag
usable.
* oacc-mem.c (acc_map_data): Adjust gomp_map_vars argument flags to
'GOMP_MAP_VARS_OPENACC | GOMP_MAP_VARS_ENTER_DATA'.
(goacc_enter_datum): Likewise for call to gomp_map_vars_async.
(goacc_enter_data_internal): Likewise.
* target.c (gomp_map_vars_existing): Add 'bool inc_ref' parameter,
conditionalize refcount increase on inc_ref. Initialize is_struct
field for tgt_var.
(gomp_map_fields_existing): Add 'bool fld_inc_ref' parameter, adjust
calls to gomp_map_vars_existing.
(gomp_map_vars_internal): Change checks of GOMP_MAP_VARS_ENTER_DATA
to use bit-and (&). Adjust OpenMP handling of GOMP_MAP_STRUCT.
(gomp_var_unref): New function, factored from code in
gomp_unmap_vars_internal.
(gomp_var_copy_back): Likewise.
(gomp_var_unref_tgt): Likewise.
(gomp_unmap_vars_internal): Reorganize unmapping logic into above
three functions, handle case when 'is_struct' is true.
(gomp_exit_data): Handle GOMP_MAP_STRUCT.
* testsuite/libgomp.c-c++-common/struct-elem-1.c: New testcase.
* testsuite/libgomp.c-c++-common/struct-elem-2.c: New testcase.
* testsuite/libgomp.c-c++-common/struct-elem-3.c: New testcase.
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 29f385c9368..4878f71ac61 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -10429,8 +10429,9 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p,
gimple_seq body, tree *list_p,
}
}
else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
- && (code == OMP_TARGET_EXIT_DATA
- || code == OACC_EXIT_DATA))
+ /* Note: we keep GOMP_MAP_STRUCT for OpenMP target exit data
+ directives, so only remove for OpenACC exit data. */
+ && code == OACC_EXIT_DATA)
remove = true;
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index da7ac037dcd..3e03f52f70a 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -964,12 +964,15 @@ struct target_var_desc {
bool always_copy_from;
/* True if this is for OpenACC 'attach'. */
bool is_attach;
+ /* True if this is a structure map. */
+ bool is_struct;
/* If GOMP_MAP_TO_PSET had a NULL pointer; used for Fortran descriptors,
which were initially unallocated. */
bool has_null_ptr_assoc;
/* Relative offset against key host_start. */
uintptr_t offset;
- /* Actual length. */
+ /* Actual length, or number of following structure elements
+ if is_struct == true. */
uintptr_t length;
};
@@ -1162,10 +1165,10 @@ struct gomp_device_descr
/* Kind of the pragma, for which gomp_map_vars () is called. */
enum gomp_map_vars_kind
{
- GOMP_MAP_VARS_OPENACC,
- GOMP_MAP_VARS_TARGET,
- GOMP_MAP_VARS_DATA,
- GOMP_MAP_VARS_ENTER_DATA
+ GOMP_MAP_VARS_OPENACC = 1,
+ GOMP_MAP_VARS_TARGET = 2,
+ GOMP_MAP_VARS_DATA = 4,
+ GOMP_MAP_VARS_ENTER_DATA = 8
};
extern void gomp_acc_declare_allocate (bool, size_t, void **, size_t *,
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 65757ab2ffc..8dc521ac6d6 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -403,7 +403,8 @@ acc_map_data (void *h, void *d, size_t s)
struct target_mem_desc *tgt
= gomp_map_vars (acc_dev, mapnum, &hostaddrs, &devaddrs, &sizes,
- &kinds, true, GOMP_MAP_VARS_ENTER_DATA);
+ &kinds, true,
+ GOMP_MAP_VARS_OPENACC | GOMP_MAP_VARS_ENTER_DATA);
assert (tgt);
assert (tgt->list_count == 1);
splay_tree_key n = tgt->list[0].key;
@@ -572,7 +573,8 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void
*kinds, int async)
struct target_mem_desc *tgt
= gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes,
- kinds, true, GOMP_MAP_VARS_ENTER_DATA);
+ kinds, true,
+ GOMP_MAP_VARS_OPENACC |
GOMP_MAP_VARS_ENTER_DATA);
assert (tgt);
assert (tgt->list_count == 1);
n = tgt->list[0].key;
@@ -1202,7 +1204,7 @@ goacc_enter_data_internal (struct gomp_device_descr
*acc_dev, size_t mapnum,
struct target_mem_desc *tgt
= gomp_map_vars_async (acc_dev, aq, groupnum, &hostaddrs[i], NULL,
&sizes[i], &kinds[i], true,
- GOMP_MAP_VARS_ENTER_DATA);
+ GOMP_MAP_VARS_OPENACC |
GOMP_MAP_VARS_ENTER_DATA);
assert (tgt);
gomp_mutex_lock (&acc_dev->lock);
diff --git a/libgomp/target.c b/libgomp/target.c
index 1a8c67c2df5..32d571337ab 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -369,7 +369,7 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
struct goacc_asyncqueue *aq, splay_tree_key oldn,
splay_tree_key newn, struct target_var_desc *tgt_var,
unsigned char kind, bool always_to_flag,
- struct gomp_coalesce_buf *cbuf)
+ bool inc_ref, struct gomp_coalesce_buf *cbuf)
{
assert (kind != GOMP_MAP_ATTACH);
@@ -377,6 +377,7 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
tgt_var->copy_from = GOMP_MAP_COPY_FROM_P (kind);
tgt_var->always_copy_from = GOMP_MAP_ALWAYS_FROM_P (kind);
tgt_var->is_attach = false;
+ tgt_var->is_struct = false;
tgt_var->offset = newn->host_start - oldn->host_start;
tgt_var->length = newn->host_end - newn->host_start;
@@ -398,7 +399,7 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
(void *) newn->host_start,
newn->host_end - newn->host_start, cbuf);
- if (oldn->refcount != REFCOUNT_INFINITY)
+ if (inc_ref && oldn->refcount != REFCOUNT_INFINITY)
oldn->refcount++;
}
@@ -453,6 +454,7 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
struct goacc_asyncqueue *aq, splay_tree_key n,
size_t first, size_t i, void **hostaddrs,
size_t *sizes, void *kinds,
+ bool fld_inc_ref,
struct gomp_coalesce_buf *cbuf)
{
struct gomp_device_descr *devicep = tgt->device_descr;
@@ -471,7 +473,7 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
&& n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset)
{
gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i],
- kind & typemask, false, cbuf);
+ kind & typemask, false, fld_inc_ref, cbuf);
return;
}
if (sizes[i] == 0)
@@ -487,7 +489,8 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
== n2->tgt_offset - n->tgt_offset)
{
gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i],
- kind & typemask, false, cbuf);
+ kind & typemask, false, fld_inc_ref,
+ cbuf);
return;
}
}
@@ -499,7 +502,7 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
&& n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset)
{
gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i],
- kind & typemask, false, cbuf);
+ kind & typemask, false, fld_inc_ref, cbuf);
return;
}
}
@@ -676,6 +679,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
size_t i, tgt_align, tgt_size, not_found_cnt = 0;
bool has_firstprivate = false;
bool has_always_ptrset = false;
+ bool openmp_p = (pragma_kind & GOMP_MAP_VARS_OPENACC) == 0;
const int rshift = short_mapkind ? 8 : 3;
const int typemask = short_mapkind ? 0xff : 0x7;
struct splay_tree_s *mem_map = &devicep->mem_map;
@@ -683,7 +687,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
struct target_mem_desc *tgt
= gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
tgt->list_count = mapnum;
- tgt->refcount = pragma_kind == GOMP_MAP_VARS_ENTER_DATA ? 0 : 1;
+ tgt->refcount = (pragma_kind & GOMP_MAP_VARS_ENTER_DATA) ? 0 : 1;
tgt->device_descr = devicep;
tgt->prev = NULL;
struct gomp_coalesce_buf cbuf, *cbufp = NULL;
@@ -798,6 +802,10 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
tgt_size = (tgt_size + align - 1) & ~(align - 1);
tgt_size += cur_node.host_end - cur_node.host_start;
not_found_cnt += last - i;
+ /* For OpenMP, we also create an entry for the struct map
+ itself, besides the elements. */
+ if (openmp_p)
+ not_found_cnt += 1;
for (i = first; i <= last; i++)
{
tgt->list[i].key = NULL;
@@ -811,9 +819,29 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
i--;
continue;
}
+
+ /* For OpenMP, structure elements do not increment refcount of the
+ splay_tree_key, only the heading struct map entry does. This is
+ to create the uniform alloc/release behavior specified in OpenMP
+ 5.0, i.e. map/unmap of just one structure element field will
+ behave the same as having alloc/release maps for all
+ (already mapped) element fields. */
+ if (openmp_p)
+ {
+ tgt->list[i].key = n;
+ tgt->list[i].offset = 0; /* Note: not OFFSET_STRUCT. */
+ tgt->list[i].copy_from = false;
+ tgt->list[i].always_copy_from = false;
+ tgt->list[i].is_attach = false;
+ tgt->list[i].is_struct = true;
+ tgt->list[i].length = sizes[i];
+
+ if (n->refcount != REFCOUNT_INFINITY)
+ n->refcount++;
+ }
for (i = first; i <= last; i++)
gomp_map_fields_existing (tgt, aq, n, first, i, hostaddrs,
- sizes, kinds, NULL);
+ sizes, kinds, !openmp_p, NULL);
i--;
continue;
}
@@ -909,7 +937,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
}
}
gomp_map_vars_existing (devicep, aq, n, &cur_node, &tgt->list[i],
- kind & typemask, always_to_cnt > 0, NULL);
+ kind & typemask, always_to_cnt > 0, true,
NULL);
i += always_to_cnt;
}
else
@@ -1064,6 +1092,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
tgt->list[j].copy_from = false;
tgt->list[j].always_copy_from = false;
tgt->list[j].is_attach = false;
+ tgt->list[j].is_struct = false;
if (k->refcount != REFCOUNT_INFINITY)
k->refcount++;
gomp_map_pointer (k->tgt, aq,
@@ -1155,11 +1184,48 @@ gomp_map_vars_internal (struct gomp_device_descr
*devicep,
field_tgt_clear = last;
tgt_size += cur_node.host_end
- (uintptr_t) hostaddrs[first];
+ if (openmp_p)
+ {
+ /* When first mapping a struct, create a region
+ encompassing all to be mapped structure elements. */
+ splay_tree_key k = &array->key;
+ k->host_start = cur_node.host_start;
+ k->host_end = cur_node.host_end;
+ k->aux = NULL;
+ k->refcount = 1;
+ k->dynamic_refcount = 0;
+ k->tgt = tgt;
+ tgt->refcount++;
+ /* Locate target address of what should be the start of
+ the entire structure. */
+ k->tgt_offset = (field_tgt_offset
+ - (hostaddrs[first] - hostaddrs[i]));
+
+ /* Save number of elements here. */
+ tgt->list[i].length = sizes[i];
+
+ tgt->list[i].copy_from = false;
+ tgt->list[i].always_copy_from = false;
+ tgt->list[i].is_attach = false;
+ tgt->list[i].is_struct = true;
+ tgt->list[i].offset = 0;
+ tgt->list[i].key = k;
+
+ array->left = NULL;
+ array->right = NULL;
+ splay_tree_insert (mem_map, array);
+ array++;
+ }
continue;
}
+
+ if (openmp_p
+ && n->refcount != REFCOUNT_INFINITY)
+ n->refcount++;
+
for (i = first; i <= last; i++)
gomp_map_fields_existing (tgt, aq, n, first, i, hostaddrs,
- sizes, kinds, cbufp);
+ sizes, kinds, !openmp_p, cbufp);
i--;
continue;
case GOMP_MAP_ALWAYS_POINTER:
@@ -1209,6 +1275,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
tgt->list[i].copy_from = false;
tgt->list[i].always_copy_from = false;
tgt->list[i].is_attach = true;
+ tgt->list[i].is_struct = false;
/* OpenACC 'attach'/'detach' doesn't affect
structured/dynamic reference counts ('n->refcount',
'n->dynamic_refcount'). */
@@ -1234,8 +1301,28 @@ gomp_map_vars_internal (struct gomp_device_descr
*devicep,
k->host_end = k->host_start + sizeof (void *);
splay_tree_key n = splay_tree_lookup (mem_map, k);
if (n && n->refcount != REFCOUNT_LINK)
- gomp_map_vars_existing (devicep, aq, n, k, &tgt->list[i],
- kind & typemask, false, cbufp);
+ {
+ if (!openmp_p)
+ gomp_map_vars_existing (devicep, aq, n, k, &tgt->list[i],
+ kind & typemask, false, true, cbufp);
+ else
+ {
+ bool inc_ref = (field_tgt_clear == FIELD_TGT_EMPTY);
+
+ gomp_map_vars_existing (devicep, aq, n, k, &tgt->list[i],
+ kind & typemask, false, inc_ref,
cbufp);
+ if (field_tgt_clear != FIELD_TGT_EMPTY)
+ {
+ k->tgt = tgt;
+ k->tgt_offset = (k->host_start
+ - field_tgt_base + field_tgt_offset);
+ if (i == field_tgt_clear)
+ field_tgt_clear = FIELD_TGT_EMPTY;
+ k->aux = NULL;
+ goto copy_map;
+ }
+ }
+ }
else
{
k->aux = NULL;
@@ -1268,6 +1355,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
tgt->list[i].always_copy_from
= GOMP_MAP_ALWAYS_FROM_P (kind & typemask);
tgt->list[i].is_attach = false;
+ tgt->list[i].is_struct = false;
tgt->list[i].offset = 0;
tgt->list[i].length = k->host_end - k->host_start;
k->refcount = 1;
@@ -1276,6 +1364,8 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
array->left = NULL;
array->right = NULL;
splay_tree_insert (mem_map, array);
+
+ copy_map:
switch (kind & typemask)
{
case GOMP_MAP_ALLOC:
@@ -1326,6 +1416,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
tgt->list[j].copy_from = false;
tgt->list[j].always_copy_from = false;
tgt->list[j].is_attach = false;
+ tgt->list[j].is_struct = false;
tgt->list[i].has_null_ptr_assoc |= !(*(void **)
hostaddrs[j]);
if (k->refcount != REFCOUNT_INFINITY)
k->refcount++;
@@ -1415,7 +1506,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
/* If the variable from "omp target enter data" map-list was already mapped,
tgt is not needed. Otherwise tgt will be freed by gomp_unmap_vars or
gomp_exit_data. */
- if (pragma_kind == GOMP_MAP_VARS_ENTER_DATA && tgt->refcount == 0)
+ if ((pragma_kind & GOMP_MAP_VARS_ENTER_DATA) && tgt->refcount == 0)
{
free (tgt);
tgt = NULL;
@@ -1523,6 +1614,50 @@ gomp_remove_var_async (struct gomp_device_descr
*devicep, splay_tree_key k,
(void) gomp_remove_var_internal (devicep, k, aq);
}
+static bool
+gomp_var_unref (splay_tree_key k)
+{
+ bool do_unmap = false;
+ if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY)
+ k->refcount--;
+ else if (k->refcount == 1)
+ {
+ k->refcount--;
+ do_unmap = true;
+ }
+ return do_unmap;
+}
+
+static void
+gomp_var_copy_back (struct target_mem_desc *tgt, bool do_unmap_and_copyfrom,
+ struct goacc_asyncqueue *aq, size_t i)
+{
+ struct gomp_device_descr *devicep = tgt->device_descr;
+ splay_tree_key k = tgt->list[i].key;
+
+ if ((do_unmap_and_copyfrom && tgt->list[i].copy_from)
+ || tgt->list[i].always_copy_from)
+ gomp_copy_dev2host (devicep, aq,
+ (void *) (k->host_start + tgt->list[i].offset),
+ (void *) (k->tgt->tgt_start + k->tgt_offset
+ + tgt->list[i].offset),
+ tgt->list[i].length);
+}
+
+static void
+gomp_var_unref_tgt (struct target_mem_desc *tgt, splay_tree_key k)
+{
+ struct gomp_device_descr *devicep = tgt->device_descr;
+ struct target_mem_desc *k_tgt = k->tgt;
+
+ bool is_tgt_unmapped = gomp_remove_var (devicep, k);
+
+ /* It would be bad if TGT got unmapped while we're still iterating
+ over its LIST_COUNT, and also expect to use it in the following
+ code. */
+ assert (!is_tgt_unmapped || k_tgt != tgt);
+}
+
/* Unmap variables described by TGT. If DO_COPYFROM is true, copy relevant
variables back from device to host: if it is false, it is assumed that this
has been done already. */
@@ -1561,42 +1696,40 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt,
bool do_copyfrom,
false, NULL);
}
- for (i = 0; i < tgt->list_count; i++)
+ for (i = 0; i < tgt->list_count;)
{
- splay_tree_key k = tgt->list[i].key;
- if (k == NULL)
- continue;
-
- /* OpenACC 'attach'/'detach' doesn't affect structured/dynamic reference
- counts ('n->refcount', 'n->dynamic_refcount'). */
- if (tgt->list[i].is_attach)
- continue;
-
- bool do_unmap = false;
- if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY)
- k->refcount--;
- else if (k->refcount == 1)
+ if (tgt->list[i].key == NULL
+ /* OpenACC 'attach'/'detach' doesn't affect structured/dynamic
+ reference counts ('n->refcount', 'n->dynamic_refcount'). */
+ || tgt->list[i].is_attach)
{
- k->refcount--;
- do_unmap = true;
+ i += 1;
+ continue;
}
- if ((do_unmap && do_copyfrom && tgt->list[i].copy_from)
- || tgt->list[i].always_copy_from)
- gomp_copy_dev2host (devicep, aq,
- (void *) (k->host_start + tgt->list[i].offset),
- (void *) (k->tgt->tgt_start + k->tgt_offset
- + tgt->list[i].offset),
- tgt->list[i].length);
- if (do_unmap)
+ if (tgt->list[i].is_struct)
+ {
+ size_t j, num_elem = tgt->list[i].length;
+
+ /* Release the struct map's reference on the splay_tree_key. */
+ bool do_unmap = gomp_var_unref (tgt->list[i].key);
+
+ for (j = i + 1; j <= i + num_elem; j++)
+ gomp_var_copy_back (tgt, do_unmap && do_copyfrom, aq, j);
+
+ if (do_unmap)
+ gomp_var_unref_tgt (tgt, tgt->list[i].key);
+ i = j;
+ }
+ else
{
- struct target_mem_desc *k_tgt = k->tgt;
- bool is_tgt_unmapped = gomp_remove_var (devicep, k);
- /* It would be bad if TGT got unmapped while we're still iterating
- over its LIST_COUNT, and also expect to use it in the following
- code. */
- assert (!is_tgt_unmapped
- || k_tgt != tgt);
+ splay_tree_key k = tgt->list[i].key;
+ bool do_unmap = gomp_var_unref (k);
+
+ gomp_var_copy_back (tgt, do_unmap && do_copyfrom, aq, i);
+ if (do_unmap)
+ gomp_var_unref_tgt (tgt, k);
+ i += 1;
}
}
@@ -2512,8 +2645,46 @@ gomp_exit_data (struct gomp_device_descr *devicep,
size_t mapnum,
cur_node.host_end - cur_node.host_start);
if (k->refcount == 0)
gomp_remove_var (devicep, k);
+ break;
+
+ case GOMP_MAP_STRUCT:
+ {
+ size_t num_elem = sizes[i];
+ cur_node.host_start = (uintptr_t) hostaddrs[i];
+ cur_node.host_end = cur_node.host_start + 1;
+
+ splay_tree_key k = splay_tree_lookup (&devicep->mem_map, &cur_node);
+ if (k)
+ {
+ for (size_t j = i + 1; j <= i + num_elem; j++)
+ {
+ cur_node.host_start = (uintptr_t) hostaddrs[j];
+ cur_node.host_end = cur_node.host_start + sizes[j];
+ splay_tree_key ek = splay_tree_lookup (&devicep->mem_map,
+ &cur_node);
+ /* All fields should lookup to same splay_tree_key. */
+ assert (ek == k);
+ }
+ bool do_unmap = gomp_var_unref (k);
+ for (size_t j = i + 1; j <= i + num_elem; j++)
+ {
+ unsigned char ekind = kinds[j] & typemask;
+ if ((ekind == GOMP_MAP_FROM && do_unmap)
+ || ekind == GOMP_MAP_ALWAYS_FROM)
+ gomp_copy_dev2host (devicep, NULL, hostaddrs[j],
+ (void *) (k->tgt->tgt_start
+ + k->tgt_offset
+ + (uintptr_t) hostaddrs[j]
+ - k->host_start), sizes[j]);
+ }
+ if (do_unmap)
+ gomp_remove_var (devicep, k);
+ }
+ i += num_elem;
+ }
break;
+
default:
gomp_mutex_unlock (&devicep->lock);
gomp_fatal ("GOMP_target_enter_exit_data unhandled kind 0x%.2x",
diff --git a/libgomp/testsuite/libgomp.c-c++-common/struct-elem-1.c
b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-1.c
new file mode 100644
index 00000000000..c49d8c12c05
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-1.c
@@ -0,0 +1,29 @@
+#include <omp.h>
+#include <stdlib.h>
+
+struct S
+{
+ int a, b;
+};
+typedef struct S S;
+
+int main (void)
+{
+ int d = omp_get_default_device ();
+ int id = omp_get_initial_device ();
+
+ if (d < 0 || d >= omp_get_num_devices ())
+ d = id;
+
+ S s;
+ #pragma omp target enter data map (alloc: s.a, s.b)
+ #pragma omp target exit data map (release: s.b)
+
+ /* OpenMP 5.0 structure element mapping rules describe that elements of same
+ structure variable should allocate/deallocate in a uniform fashion, so
+ "s.a" should be removed together by above 'exit data'. */
+ if (omp_target_is_present (&s.a, d))
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/struct-elem-2.c
b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-2.c
new file mode 100644
index 00000000000..555c6e3e8e0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-2.c
@@ -0,0 +1,44 @@
+#include <omp.h>
+#include <stdlib.h>
+
+struct S
+{
+ int a, b, c, d;
+};
+typedef struct S S;
+
+int main (void)
+{
+ int d = omp_get_default_device ();
+ int id = omp_get_initial_device ();
+
+ if (d < 0 || d >= omp_get_num_devices ())
+ d = id;
+
+ S s;
+ #pragma omp target enter data map (alloc: s.a, s.b, s.c, s.d)
+ #pragma omp target enter data map (alloc: s.c)
+ #pragma omp target enter data map (alloc: s.b, s.d)
+ #pragma omp target enter data map (alloc: s.a, s.c, s.b)
+
+ #pragma omp target exit data map (release: s.a)
+ #pragma omp target exit data map (release: s.d)
+ #pragma omp target exit data map (release: s.c)
+ #pragma omp target exit data map (release: s.b)
+
+ /* OpenMP 5.0 structure element mapping rules describe that elements of same
+ structure variable should allocate/deallocate in a uniform fashion, so
+ all elements of 's' should be removed together by above 'exit data's. */
+ if (omp_target_is_present (&s, d))
+ abort ();
+ if (omp_target_is_present (&s.a, d))
+ abort ();
+ if (omp_target_is_present (&s.b, d))
+ abort ();
+ if (omp_target_is_present (&s.c, d))
+ abort ();
+ if (omp_target_is_present (&s.d, d))
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/struct-elem-3.c
b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-3.c
new file mode 100644
index 00000000000..4850eabd879
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/struct-elem-3.c
@@ -0,0 +1,63 @@
+#include <omp.h>
+#include <stdlib.h>
+
+struct S
+{
+ int a, b, c, d;
+};
+typedef struct S S;
+
+int main (void)
+{
+ int d = omp_get_default_device ();
+ int id = omp_get_initial_device ();
+
+ if (d < 0 || d >= omp_get_num_devices ())
+ d = id;
+
+ S s;
+
+ #pragma omp target enter data map (alloc: s)
+ #pragma omp target enter data map (alloc: s)
+
+ #pragma omp target exit data map (release: s.a)
+ #pragma omp target exit data map (release: s.b)
+
+ /* OpenMP 5.0 structure element mapping rules describe that elements of same
+ structure variable should allocate/deallocate in a uniform fashion, so
+ all elements of 's' should be removed together by above 'exit data's. */
+ if (omp_target_is_present (&s, d))
+ abort ();
+ if (omp_target_is_present (&s.a, d))
+ abort ();
+ if (omp_target_is_present (&s.b, d))
+ abort ();
+ if (omp_target_is_present (&s.c, d))
+ abort ();
+ if (omp_target_is_present (&s.d, d))
+ abort ();
+
+ #pragma omp target enter data map (alloc: s.a, s.b)
+ #pragma omp target enter data map (alloc: s.a)
+ #pragma omp target enter data map (alloc: s.b)
+
+ #pragma omp target exit data map (release: s)
+ #pragma omp target exit data map (release: s)
+ #pragma omp target exit data map (release: s)
+
+ /* OpenMP 5.0 structure element mapping rules describe that elements of same
+ structure variable should allocate/deallocate in a uniform fashion, so
+ all elements of 's' should be removed together by above 'exit data's. */
+ if (omp_target_is_present (&s, d))
+ abort ();
+ if (omp_target_is_present (&s.a, d))
+ abort ();
+ if (omp_target_is_present (&s.b, d))
+ abort ();
+ if (omp_target_is_present (&s.c, d))
+ abort ();
+ if (omp_target_is_present (&s.d, d))
+ abort ();
+
+ return 0;
+}
