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