This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[gomp4.1] Start of structure element mapping support
- 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: Fri, 31 Jul 2015 18:16:10 +0200
- Subject: [gomp4.1] Start of structure element mapping support
- Authentication-results: sourceware.org; auth=none
- Reply-to: Jakub Jelinek <jakub at redhat dot com>
Hi!
This patch is the start of implementation of struct element mapping.
I'm not handling structure element based array sections (neither
array based, nor pointer/reference based) yet, nor C++.
If the whole struct is already mapped, then that mapping is used,
otherwise we require that either all the fields are already mapped, or none
of them (otherwise runtime error). If none of them, then we allocate
enough room for the first to last mapped field, and place all the individual
allocations into the allocated space.
2015-07-31 Jakub Jelinek <jakub@redhat.com>
* gimplify.c (gimplify_scan_omp_clauses): Handle
map clauses with COMPONENT_REF.
* omp-low.c (lower_omp_target): Handle GOMP_MAP_STRUCT.
Handle GOMP_MAP_RELEASE for zero-length array sections.
* tree-pretty-print.c (dump_omp_clause): Handle
GOMP_MAP_STRUCT.
gcc/c/
* c-parser.c (c_parser_omp_variable_list): Parse struct
elements in map/to/from clauses.
* c-typeck.c (handle_omp_array_sections): Handle
GOMP_MAP_RELEASE for zero-length array sections.
(c_finish_omp_clauses): Handle struct elements in
map/to/from OpenMP clauses.
gcc/cp/
* semantics.c (handle_omp_array_sections): Handle
GOMP_MAP_RELEASE for zero-length array sections.
include/
* gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_STRUCT.
libgomp/
* target.c (gomp_map_fields_existing): New function.
(gomp_map_vars): Handle GOMP_MAP_STRUCT.
* testsuite/libgomp.c/target-21.c: New test.
--- gcc/gimplify.c.jj 2015-07-31 16:55:01.482411392 +0200
+++ gcc/gimplify.c 2015-07-31 16:57:22.307320290 +0200
@@ -6202,6 +6202,7 @@ gimplify_scan_omp_clauses (tree *list_p,
{
struct gimplify_omp_ctx *ctx, *outer_ctx;
tree c;
+ hash_map<tree, tree> *struct_map_to_clause = NULL;
ctx = new_omp_context (region_type);
outer_ctx = ctx->outer_context;
@@ -6442,6 +6443,11 @@ gimplify_scan_omp_clauses (tree *list_p,
}
if (!DECL_P (decl))
{
+ if (TREE_CODE (decl) == COMPONENT_REF)
+ {
+ while (TREE_CODE (decl) == COMPONENT_REF)
+ decl = TREE_OPERAND (decl, 0);
+ }
if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p,
NULL, is_gimple_lvalue, fb_lvalue)
== GS_ERROR)
@@ -6449,6 +6455,128 @@ gimplify_scan_omp_clauses (tree *list_p,
remove = true;
break;
}
+ if (DECL_P (decl))
+ {
+ if (error_operand_p (decl))
+ {
+ remove = true;
+ break;
+ }
+
+ if (TYPE_SIZE_UNIT (TREE_TYPE (decl)) == NULL
+ || (TREE_CODE (TYPE_SIZE_UNIT (TREE_TYPE (decl)))
+ != INTEGER_CST))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "mapping field %qE of variable length "
+ "structure", OMP_CLAUSE_DECL (c));
+ remove = true;
+ break;
+ }
+
+ tree offset;
+ HOST_WIDE_INT bitsize, bitpos;
+ machine_mode mode;
+ int unsignedp, volatilep = 0;
+ tree base
+ = get_inner_reference (OMP_CLAUSE_DECL (c), &bitsize,
+ &bitpos, &offset, &mode, &unsignedp,
+ &volatilep, false);
+ gcc_assert (base == decl
+ && (offset == NULL_TREE
+ || TREE_CODE (offset) == INTEGER_CST));
+
+ splay_tree_node n
+ = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
+ if (n == NULL || (n->value & GOVD_MAP) == 0)
+ {
+ *list_p = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (*list_p, GOMP_MAP_STRUCT);
+ OMP_CLAUSE_DECL (*list_p) = decl;
+ OMP_CLAUSE_SIZE (*list_p) = size_int (1);
+ OMP_CLAUSE_CHAIN (*list_p) = c;
+ if (struct_map_to_clause == NULL)
+ struct_map_to_clause = new hash_map<tree, tree>;
+ struct_map_to_clause->put (decl, *list_p);
+ list_p = &OMP_CLAUSE_CHAIN (*list_p);
+ flags = GOVD_MAP | GOVD_EXPLICIT;
+ if (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS)
+ flags |= GOVD_SEEN;
+ goto do_add_decl;
+ }
+ else
+ {
+ tree *osc = struct_map_to_clause->get (decl), *sc;
+ if (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS)
+ n->value |= GOVD_SEEN;
+ offset_int o1, o2;
+ if (offset)
+ o1 = wi::to_offset (offset);
+ else
+ o1 = 0;
+ if (bitpos)
+ o1 = o1 + bitpos / BITS_PER_UNIT;
+ for (sc = &OMP_CLAUSE_CHAIN (*osc); *sc != c;
+ sc = &OMP_CLAUSE_CHAIN (*sc))
+ if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) != COMPONENT_REF)
+ break;
+ else
+ {
+ tree offset2;
+ HOST_WIDE_INT bitsize2, bitpos2;
+ base = get_inner_reference (OMP_CLAUSE_DECL (*sc),
+ &bitsize2, &bitpos2,
+ &offset2, &mode,
+ &unsignedp, &volatilep,
+ false);
+ if (base != decl)
+ break;
+ gcc_assert (offset == NULL_TREE
+ || TREE_CODE (offset) == INTEGER_CST);
+ tree d1 = OMP_CLAUSE_DECL (*sc);
+ tree d2 = OMP_CLAUSE_DECL (c);
+ while (TREE_CODE (d1) == COMPONENT_REF)
+ if (TREE_CODE (d2) == COMPONENT_REF
+ && TREE_OPERAND (d1, 1)
+ == TREE_OPERAND (d2, 1))
+ {
+ d1 = TREE_OPERAND (d1, 0);
+ d2 = TREE_OPERAND (d2, 0);
+ }
+ else
+ break;
+ if (d1 == d2)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qE appears more than once in map "
+ "clauses", OMP_CLAUSE_DECL (c));
+ remove = true;
+ break;
+ }
+ if (offset2)
+ o2 = wi::to_offset (offset2);
+ else
+ o2 = 0;
+ if (bitpos2)
+ o2 = o2 + bitpos2 / BITS_PER_UNIT;
+ if (wi::ltu_p (o1, o2)
+ || (wi::eq_p (o1, o2) && bitpos < bitpos2))
+ break;
+ }
+ if (!remove)
+ OMP_CLAUSE_SIZE (*osc)
+ = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
+ size_one_node);
+ if (!remove && *sc != c)
+ {
+ *list_p = OMP_CLAUSE_CHAIN (c);
+ OMP_CLAUSE_CHAIN (c) = *sc;
+ *sc = c;
+ continue;
+ }
+ }
+ }
break;
}
flags = GOVD_MAP | GOVD_EXPLICIT;
@@ -6790,6 +6918,8 @@ gimplify_scan_omp_clauses (tree *list_p,
}
gimplify_omp_ctxp = ctx;
+ if (struct_map_to_clause)
+ delete struct_map_to_clause;
}
struct gimplify_adjust_omp_clauses_data
--- gcc/omp-low.c.jj 2015-07-31 16:55:01.272414510 +0200
+++ gcc/omp-low.c 2015-07-31 16:57:22.317320141 +0200
@@ -12954,6 +12954,7 @@ lower_omp_target (gimple_stmt_iterator *
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_STRUCT:
break;
case GOMP_MAP_FORCE_ALLOC:
case GOMP_MAP_FORCE_TO:
@@ -13303,6 +13304,7 @@ lower_omp_target (gimple_stmt_iterator *
case GOMP_MAP_ALWAYS_TO:
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_RELEASE:
tkind_zero = GOMP_MAP_ZERO_LEN_ARRAY_SECTION;
break;
default:
--- gcc/tree-pretty-print.c.jj 2015-07-31 16:55:01.484411362 +0200
+++ gcc/tree-pretty-print.c 2015-07-31 16:57:22.320320097 +0200
@@ -643,6 +643,9 @@ dump_omp_clause (pretty_printer *pp, tre
case GOMP_MAP_FIRSTPRIVATE_POINTER:
pp_string (pp, "firstprivate");
break;
+ case GOMP_MAP_STRUCT:
+ pp_string (pp, "struct");
+ break;
default:
gcc_unreachable ();
}
--- gcc/c/c-parser.c.jj 2015-07-31 16:55:01.481411407 +0200
+++ gcc/c/c-parser.c 2015-07-31 16:57:22.313320201 +0200
@@ -10190,10 +10190,25 @@ c_parser_omp_variable_list (c_parser *pa
t = error_mark_node;
break;
}
- /* FALL THROUGH. */
+ /* FALLTHROUGH */
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_FROM:
case OMP_CLAUSE_TO:
+ while (c_parser_next_token_is (parser, CPP_DOT))
+ {
+ location_t op_loc = c_parser_peek_token (parser)->location;
+ c_parser_consume_token (parser);
+ if (!c_parser_next_token_is (parser, CPP_NAME))
+ {
+ c_parser_error (parser, "expected identifier");
+ t = error_mark_node;
+ break;
+ }
+ tree ident = c_parser_peek_token (parser)->value;
+ c_parser_consume_token (parser);
+ t = build_component_ref (op_loc, t, ident);
+ }
+ /* FALLTHROUGH */
case OMP_CLAUSE_DEPEND:
case OMP_CLAUSE_REDUCTION:
while (c_parser_next_token_is (parser, CPP_OPEN_SQUARE))
--- gcc/c/c-typeck.c.jj 2015-07-31 16:55:01.482411392 +0200
+++ gcc/c/c-typeck.c 2015-07-31 16:58:09.246623290 +0200
@@ -12040,6 +12040,7 @@ handle_omp_array_sections (tree c, bool
case GOMP_MAP_ALWAYS_TO:
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_RELEASE:
OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
break;
default:
@@ -12117,7 +12118,7 @@ tree
c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd)
{
bitmap_head generic_head, firstprivate_head, lastprivate_head;
- bitmap_head aligned_head, map_head;
+ bitmap_head aligned_head, map_head, map_field_head;
tree c, t, type, *pc;
tree simdlen = NULL_TREE, safelen = NULL_TREE;
bool branch_seen = false;
@@ -12130,6 +12131,7 @@ c_finish_omp_clauses (tree clauses, bool
bitmap_initialize (&lastprivate_head, &bitmap_default_obstack);
bitmap_initialize (&aligned_head, &bitmap_default_obstack);
bitmap_initialize (&map_head, &bitmap_default_obstack);
+ bitmap_initialize (&map_field_head, &bitmap_default_obstack);
for (pc = &clauses, c = clauses; c ; c = *pc)
{
@@ -12574,8 +12576,49 @@ c_finish_omp_clauses (tree clauses, bool
break;
}
if (t == error_mark_node)
- remove = true;
- else if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
+ {
+ remove = true;
+ break;
+ }
+ if (TREE_CODE (t) == COMPONENT_REF
+ && is_omp
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
+ {
+ if (DECL_BIT_FIELD (TREE_OPERAND (t, 1)))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "bit-field %qE in %qs clause",
+ t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ remove = true;
+ }
+ else if (!lang_hooks.types.omp_mappable_type (TREE_TYPE (t)))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qE does not have a mappable type in %qs clause",
+ t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ remove = true;
+ }
+ while (TREE_CODE (t) == COMPONENT_REF)
+ {
+ if (TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0)))
+ == UNION_TYPE)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qE is a member of a union", t);
+ remove = true;
+ break;
+ }
+ t = TREE_OPERAND (t, 0);
+ }
+ if (remove)
+ break;
+ if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
+ {
+ if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ break;
+ }
+ }
+ if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qE is not a variable in %qs clause", t,
@@ -12597,6 +12640,7 @@ c_finish_omp_clauses (tree clauses, bool
== GOMP_MAP_FIRSTPRIVATE_POINTER)
|| (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_FORCE_DEVICEPTR)))
+ && t == OMP_CLAUSE_DECL (c)
&& !lang_hooks.types.omp_mappable_type (TREE_TYPE (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
@@ -12613,7 +12657,12 @@ c_finish_omp_clauses (tree clauses, bool
remove = true;
}
else
- bitmap_set_bit (&map_head, DECL_UID (t));
+ {
+ bitmap_set_bit (&map_head, DECL_UID (t));
+ if (t != OMP_CLAUSE_DECL (c)
+ && TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF)
+ bitmap_set_bit (&map_field_head, DECL_UID (t));
+ }
break;
case OMP_CLAUSE_TO_DECLARE:
--- gcc/cp/semantics.c.jj 2015-07-31 16:55:01.485411348 +0200
+++ gcc/cp/semantics.c 2015-07-31 16:57:22.303320349 +0200
@@ -4836,6 +4836,7 @@ handle_omp_array_sections (tree c, bool
case GOMP_MAP_ALWAYS_TO:
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_RELEASE:
OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
break;
default:
--- include/gomp-constants.h.jj 2015-07-31 16:55:01.604409581 +0200
+++ include/gomp-constants.h 2015-07-31 16:55:38.711858574 +0200
@@ -102,6 +102,14 @@ enum gomp_map_kind
/* If not already present, allocate. And unconditionally copy to and from
device. */
GOMP_MAP_ALWAYS_TOFROM = (GOMP_MAP_FLAG_ALWAYS | GOMP_MAP_TOFROM),
+ /* Map a sparse struct; the address is the base of the structure, alignment
+ it's required alignment, and size is the number of adjacent entries
+ that belong to the struct. The adjacent entries should be sorted by
+ increasing address, so it is easy to determine lowest needed address
+ (address of the first adjacent entry) and highest needed address
+ (address of the last adjacent entry plus its size). */
+ GOMP_MAP_STRUCT = (GOMP_MAP_FLAG_ALWAYS
+ | GOMP_MAP_FLAG_SPECIAL | 0),
/* OpenMP 4.1 alias for forced deallocation. */
GOMP_MAP_DELETE = GOMP_MAP_FORCE_DEALLOC,
/* Decrement usage count and deallocate if zero. */
--- libgomp/target.c.jj 2015-07-31 16:55:01.981403983 +0200
+++ libgomp/target.c 2015-07-31 16:55:38.710858589 +0200
@@ -245,6 +245,66 @@ gomp_map_pointer (struct target_mem_desc
sizeof (void *));
}
+static void
+gomp_map_fields_existing (struct target_mem_desc *tgt, splay_tree_key n,
+ size_t first, size_t i, void **hostaddrs,
+ size_t *sizes, void *kinds)
+{
+ struct gomp_device_descr *devicep = tgt->device_descr;
+ struct splay_tree_s *mem_map = &devicep->mem_map;
+ struct splay_tree_key_s cur_node;
+ int kind;
+ const bool short_mapkind = true;
+ const int typemask = short_mapkind ? 0xff : 0x7;
+
+ cur_node.host_start = (uintptr_t) hostaddrs[i];
+ cur_node.host_end = cur_node.host_start + sizes[i];
+ splay_tree_key n2 = splay_tree_lookup (mem_map, &cur_node);
+ kind = get_kind (short_mapkind, kinds, i);
+ if (n2
+ && n2->tgt == n->tgt
+ && n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset)
+ {
+ gomp_map_vars_existing (devicep, n2, &cur_node,
+ &tgt->list[i], kind & typemask);
+ return;
+ }
+ if (sizes[i] == 0)
+ {
+ if (cur_node.host_start > (uintptr_t) hostaddrs[first - 1])
+ {
+ cur_node.host_start--;
+ n2 = splay_tree_lookup (mem_map, &cur_node);
+ cur_node.host_start++;
+ if (n2
+ && n2->tgt == n->tgt
+ && n2->host_start - n->host_start
+ == n2->tgt_offset - n->tgt_offset)
+ {
+ gomp_map_vars_existing (devicep, n2, &cur_node, &tgt->list[i],
+ kind & typemask);
+ return;
+ }
+ }
+ cur_node.host_end++;
+ n2 = splay_tree_lookup (mem_map, &cur_node);
+ cur_node.host_end--;
+ if (n2
+ && n2->tgt == n->tgt
+ && n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset)
+ {
+ gomp_map_vars_existing (devicep, n2, &cur_node, &tgt->list[i],
+ kind & typemask);
+ return;
+ }
+ }
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("Trying to map into device [%p..%p) structure element when "
+ "other mapped elements from the same structure weren't mapped "
+ "together with it", (void *) cur_node.host_start,
+ (void *) cur_node.host_end);
+}
+
attribute_hidden struct target_mem_desc *
gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds,
@@ -304,6 +364,37 @@ gomp_map_vars (struct gomp_device_descr
tgt->list[i].offset = ~(uintptr_t) 0;
continue;
}
+ else if ((kind & typemask) == GOMP_MAP_STRUCT)
+ {
+ size_t first = i + 1;
+ size_t last = i + sizes[i];
+ cur_node.host_start = (uintptr_t) hostaddrs[i];
+ cur_node.host_end = (uintptr_t) hostaddrs[last]
+ + sizes[last];
+ tgt->list[i].key = NULL;
+ tgt->list[i].offset = ~(uintptr_t) 2;
+ splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
+ if (n == NULL)
+ {
+ size_t align = (size_t) 1 << (kind >> rshift);
+ if (tgt_align < align)
+ tgt_align = align;
+ tgt_size -= (uintptr_t) hostaddrs[first]
+ - (uintptr_t) hostaddrs[i];
+ tgt_size = (tgt_size + align - 1) & ~(align - 1);
+ tgt_size += cur_node.host_end - (uintptr_t) hostaddrs[i];
+ not_found_cnt += last - i;
+ for (i = first; i <= last; i++)
+ tgt->list[i].key = NULL;
+ i--;
+ continue;
+ }
+ for (i = first; i <= last; i++)
+ gomp_map_fields_existing (tgt, n, first, i, hostaddrs,
+ sizes, kinds);
+ i--;
+ continue;
+ }
cur_node.host_start = (uintptr_t) hostaddrs[i];
if (!GOMP_MAP_POINTER_P (kind & typemask))
cur_node.host_end = cur_node.host_start + sizes[i];
@@ -406,7 +497,8 @@ gomp_map_vars (struct gomp_device_descr
if (not_found_cnt)
tgt->array = gomp_malloc (not_found_cnt * sizeof (*tgt->array));
splay_tree_node array = tgt->array;
- size_t j;
+ size_t j, field_tgt_offset = 0, field_tgt_clear = ~(size_t) 0;
+ uintptr_t field_tgt_base = 0;
for (i = 0; i < mapnum; i++)
if (tgt->list[i].key == NULL)
@@ -414,24 +506,53 @@ gomp_map_vars (struct gomp_device_descr
int kind = get_kind (short_mapkind, kinds, i);
if (hostaddrs[i] == NULL)
continue;
- if ((kind & typemask) == GOMP_MAP_FIRSTPRIVATE)
+ switch (kind & typemask)
{
- size_t align = (size_t) 1 << (kind >> rshift);
+ size_t align, len, first, last;
+ splay_tree_key n;
+ case GOMP_MAP_FIRSTPRIVATE:
+ align = (size_t) 1 << (kind >> rshift);
tgt_size = (tgt_size + align - 1) & ~(align - 1);
tgt->list[i].offset = tgt_size;
- size_t len = sizes[i];
+ len = sizes[i];
devicep->host2dev_func (devicep->target_id,
(void *) (tgt->tgt_start + tgt_size),
(void *) hostaddrs[i], len);
tgt_size += len;
continue;
- }
- switch (kind & typemask)
- {
case GOMP_MAP_FIRSTPRIVATE_INT:
case GOMP_MAP_USE_DEVICE_PTR:
case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
continue;
+ case GOMP_MAP_STRUCT:
+ first = i + 1;
+ last = i + sizes[i];
+ cur_node.host_start = (uintptr_t) hostaddrs[i];
+ cur_node.host_end = (uintptr_t) hostaddrs[last]
+ + sizes[last];
+ if (tgt->list[first].key != NULL)
+ continue;
+ n = splay_tree_lookup (mem_map, &cur_node);
+ if (n == NULL)
+ {
+ size_t align = (size_t) 1 << (kind >> rshift);
+ tgt_size -= (uintptr_t) hostaddrs[first]
+ - (uintptr_t) hostaddrs[i];
+ tgt_size = (tgt_size + align - 1) & ~(align - 1);
+ tgt_size += (uintptr_t) hostaddrs[first]
+ - (uintptr_t) hostaddrs[i];
+ field_tgt_base = (uintptr_t) hostaddrs[first];
+ field_tgt_offset = tgt_size;
+ field_tgt_clear = last;
+ tgt_size += cur_node.host_end
+ - (uintptr_t) hostaddrs[first];
+ continue;
+ }
+ for (i = first; i <= last; i++)
+ gomp_map_fields_existing (tgt, n, first, i, hostaddrs,
+ sizes, kinds);
+ i--;
+ continue;
default:
break;
}
@@ -449,10 +570,20 @@ gomp_map_vars (struct gomp_device_descr
{
size_t align = (size_t) 1 << (kind >> rshift);
tgt->list[i].key = k;
- tgt_size = (tgt_size + align - 1) & ~(align - 1);
k->tgt = tgt;
- k->tgt_offset = tgt_size;
- tgt_size += k->host_end - k->host_start;
+ if (field_tgt_clear != ~(size_t) 0)
+ {
+ k->tgt_offset = k->host_start - field_tgt_base
+ + field_tgt_offset;
+ if (i == field_tgt_clear)
+ field_tgt_clear = ~(size_t) 0;
+ }
+ else
+ {
+ tgt_size = (tgt_size + align - 1) & ~(align - 1);
+ k->tgt_offset = tgt_size;
+ tgt_size += k->host_end - k->host_start;
+ }
tgt->list[i].copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask);
tgt->list[i].always_copy_from
= GOMP_MAP_ALWAYS_FROM_P (kind & typemask);
@@ -571,6 +702,12 @@ gomp_map_vars (struct gomp_device_descr
cur_node.tgt_offset = (uintptr_t) hostaddrs[i];
else if (tgt->list[i].offset == ~(uintptr_t) 1)
cur_node.tgt_offset = 0;
+ else if (tgt->list[i].offset == ~(uintptr_t) 2)
+ cur_node.tgt_offset = tgt->list[i + 1].key->tgt->tgt_start
+ + tgt->list[i + 1].key->tgt_offset
+ + tgt->list[i + 1].offset
+ + (uintptr_t) hostaddrs[i]
+ - (uintptr_t) hostaddrs[i + 1];
else
cur_node.tgt_offset = tgt->tgt_start
+ tgt->list[i].offset;
--- libgomp/testsuite/libgomp.c/target-21.c.jj 2015-07-31 17:00:30.415527080 +0200
+++ libgomp/testsuite/libgomp.c/target-21.c 2015-07-31 17:32:56.098638516 +0200
@@ -0,0 +1,55 @@
+extern void abort (void);
+union U { int x; long long y; };
+struct T { int a; union U b; int c; };
+struct S { int s; int u; struct T v; union U w; };
+
+int
+main ()
+{
+ struct S s;
+ s.s = 0;
+ s.u = 1;
+ s.v.a = 2;
+ s.v.b.y = 3LL;
+ s.v.c = 19;
+ s.w.x = 4;
+ int err = 0;
+ #pragma omp target map (to:s.v.b, s.u) map (from: s.w, err)
+ {
+ err = 0;
+ if (s.u != 1 || s.v.b.y != 3LL)
+ err = 1;
+ s.w.x = 6;
+ }
+ if (err || s.w.x != 6)
+ abort ();
+ s.u++;
+ s.v.a++;
+ s.v.b.y++;
+ s.w.x++;
+ #pragma omp target data map (tofrom: s)
+ #pragma omp target map (always to: s.w, err) map (alloc:s.u, s.v.b)
+ {
+ err = 0;
+ if (s.u != 2 || s.v.b.y != 4LL || s.w.x != 7)
+ err = 1;
+ s.w.x = 8;
+ }
+ if (err || s.w.x != 8)
+ abort ();
+ s.u++;
+ s.v.a++;
+ s.v.b.y++;
+ s.w.x++;
+ #pragma omp target data map (from: s.w) map (to: s.v.b, s.u)
+ #pragma omp target map (always to: s.w, err) map (alloc:s.u, s.v.b)
+ {
+ err = 0;
+ if (s.u != 3 || s.v.b.y != 5LL || s.w.x != 9)
+ err = 1;
+ s.w.x = 11;
+ }
+ if (err || s.w.x != 11)
+ abort ();
+ return 0;
+}
Jakub