* [gomp4.1] Start of structure element mapping support
@ 2015-07-31 16:28 Jakub Jelinek
2015-08-28 18:23 ` [gomp4.1] WIP: Structure " Jakub Jelinek
2019-10-16 13:35 ` [gomp4.1] Start of structure element mapping support Thomas Schwinge
0 siblings, 2 replies; 11+ messages in thread
From: Jakub Jelinek @ 2015-07-31 16:28 UTC (permalink / raw)
To: Ilya Verbin; +Cc: gcc-patches
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
^ permalink raw reply [flat|nested] 11+ messages in thread
* [gomp4.1] WIP: Structure element mapping support
2015-07-31 16:28 [gomp4.1] Start of structure element mapping support Jakub Jelinek
@ 2015-08-28 18:23 ` Jakub Jelinek
2015-08-31 15:08 ` [gomp4.1] " Jakub Jelinek
2019-10-16 13:35 ` [gomp4.1] Start of structure element mapping support Thomas Schwinge
1 sibling, 1 reply; 11+ messages in thread
From: Jakub Jelinek @ 2015-08-28 18:23 UTC (permalink / raw)
To: Ilya Verbin; +Cc: gcc-patches
Hi!
Here is my current WIP on further structure element mapping support
(so, structure element {pointer,reference to pointer,reference to array}
based array sections, start of C++ support (still need to add tests for
template instantiation and verify it works properly)).
I have still pending questions on mapping of references (other than
array sections) and structure element references pending, hope they will be
responded to soon and will be able to commit this next week.
--- gcc/gimplify.c.jj 2015-08-24 14:32:06.000000000 +0200
+++ gcc/gimplify.c 2015-08-28 19:18:15.551860807 +0200
@@ -6203,6 +6203,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;
+ tree *orig_list_p = list_p;
ctx = new_omp_context (region_type);
outer_ctx = ctx->outer_context;
@@ -6443,13 +6444,31 @@ gimplify_scan_omp_clauses (tree *list_p,
}
if (!DECL_P (decl))
{
+ tree d = decl, *pd;
+ if (TREE_CODE (d) == ARRAY_REF)
+ {
+ while (TREE_CODE (d) == ARRAY_REF)
+ d = TREE_OPERAND (d, 0);
+ if (TREE_CODE (d) == COMPONENT_REF
+ && TREE_CODE (TREE_TYPE (d)) == ARRAY_TYPE)
+ decl = d;
+ }
+ pd = &OMP_CLAUSE_DECL (c);
+ if (d == decl
+ && TREE_CODE (decl) == INDIRECT_REF
+ && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+ == REFERENCE_TYPE))
+ {
+ pd = &TREE_OPERAND (decl, 0);
+ decl = TREE_OPERAND (decl, 0);
+ }
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)
+ if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
== GS_ERROR)
{
remove = true;
@@ -6478,18 +6497,48 @@ gimplify_scan_omp_clauses (tree *list_p,
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);
+ tree base = OMP_CLAUSE_DECL (c);
+ while (TREE_CODE (base) == ARRAY_REF)
+ base = TREE_OPERAND (base, 0);
+ if (TREE_CODE (base) == INDIRECT_REF)
+ base = TREE_OPERAND (base, 0);
+ base = get_inner_reference (base, &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)
+ bool ptr = (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_POINTER);
+ if (n == NULL || (n->value & (ptr ? GOVD_PRIVATE
+ : GOVD_MAP)) == 0)
{
+ if (ptr)
+ {
+ tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_PRIVATE);
+ OMP_CLAUSE_DECL (c2) = decl;
+ OMP_CLAUSE_CHAIN (c2) = *orig_list_p;
+ *orig_list_p = c2;
+ if (struct_map_to_clause == NULL)
+ struct_map_to_clause = new hash_map<tree, tree>;
+ tree *osc;
+ if (n == NULL || (n->value & GOVD_MAP) == 0)
+ osc = NULL;
+ else
+ osc = struct_map_to_clause->get (decl);
+ if (osc == NULL)
+ struct_map_to_clause->put (decl,
+ tree_cons (NULL_TREE,
+ c, NULL_TREE));
+ else
+ *osc = tree_cons (*osc, c, NULL_TREE);
+ flags = GOVD_PRIVATE | GOVD_EXPLICIT;
+ goto do_add_decl;
+ }
*list_p = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (*list_p, GOMP_MAP_STRUCT);
@@ -6508,6 +6557,9 @@ gimplify_scan_omp_clauses (tree *list_p,
else
{
tree *osc = struct_map_to_clause->get (decl), *sc;
+ tree *pt = NULL;
+ if (!ptr && TREE_CODE (*osc) == TREE_LIST)
+ osc = &TREE_PURPOSE (*osc);
if (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS)
n->value |= GOVD_SEEN;
offset_int o1, o2;
@@ -6517,25 +6569,58 @@ gimplify_scan_omp_clauses (tree *list_p,
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)
+ if (ptr)
+ pt = osc;
+ else
+ sc = &OMP_CLAUSE_CHAIN (*osc);
+ for (; ptr ? (*pt && (sc = &TREE_VALUE (*pt)))
+ : *sc != c;
+ ptr ? (pt = &TREE_CHAIN (*pt))
+ : (sc = &OMP_CLAUSE_CHAIN (*sc)))
+ if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) != COMPONENT_REF
+ && (TREE_CODE (OMP_CLAUSE_DECL (*sc))
+ != INDIRECT_REF)
+ && TREE_CODE (OMP_CLAUSE_DECL (*sc)) != ARRAY_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);
+ base = OMP_CLAUSE_DECL (*sc);
+ if (TREE_CODE (base) == ARRAY_REF)
+ {
+ while (TREE_CODE (base) == ARRAY_REF)
+ base = TREE_OPERAND (base, 0);
+ if (TREE_CODE (base) != COMPONENT_REF
+ || (TREE_CODE (TREE_TYPE (base))
+ != ARRAY_TYPE))
+ break;
+ }
+ else if (TREE_CODE (base) == INDIRECT_REF
+ && (TREE_CODE (TREE_OPERAND (base, 0))
+ == COMPONENT_REF)
+ && (TREE_CODE (TREE_TYPE
+ (TREE_OPERAND (base, 0)))
+ == REFERENCE_TYPE))
+ base = TREE_OPERAND (base, 0);
+ base = get_inner_reference (base, &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) == ARRAY_REF)
+ d1 = TREE_OPERAND (d1, 0);
+ while (TREE_CODE (d2) == ARRAY_REF)
+ d2 = TREE_OPERAND (d2, 0);
+ if (TREE_CODE (d1) == INDIRECT_REF)
+ d1 = TREE_OPERAND (d1, 0);
+ if (TREE_CODE (d2) == INDIRECT_REF)
+ d2 = TREE_OPERAND (d2, 0);
while (TREE_CODE (d1) == COMPONENT_REF)
if (TREE_CODE (d2) == COMPONENT_REF
&& TREE_OPERAND (d1, 1)
@@ -6564,6 +6649,12 @@ gimplify_scan_omp_clauses (tree *list_p,
|| (wi::eq_p (o1, o2) && bitpos < bitpos2))
break;
}
+ if (ptr)
+ {
+ if (!remove)
+ *pt = tree_cons (TREE_PURPOSE (*osc), c, *pt);
+ break;
+ }
if (!remove)
OMP_CLAUSE_SIZE (*osc)
= size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
@@ -7176,12 +7267,48 @@ gimplify_adjust_omp_clauses (gimple_seq
case OMP_CLAUSE_MAP:
decl = OMP_CLAUSE_DECL (c);
if (!DECL_P (decl))
- break;
+ {
+ if ((ctx->region_type & ORT_TARGET) != 0
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ {
+ if (TREE_CODE (decl) == INDIRECT_REF
+ && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+ == REFERENCE_TYPE))
+ decl = TREE_OPERAND (decl, 0);
+ if (TREE_CODE (decl) == COMPONENT_REF)
+ {
+ while (TREE_CODE (decl) == COMPONENT_REF)
+ decl = TREE_OPERAND (decl, 0);
+ if (DECL_P (decl))
+ {
+ n = splay_tree_lookup (ctx->variables,
+ (splay_tree_key) decl);
+ if (!(n->value & GOVD_SEEN))
+ remove = true;
+ }
+ }
+ }
+ break;
+ }
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
if ((ctx->region_type & ORT_TARGET) != 0
&& !(n->value & GOVD_SEEN)
- && !(OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS))
- remove = true;
+ && ((OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS) == 0
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT))
+ {
+ remove = true;
+ /* For struct element mapping, if struct is never referenced
+ in target block and none of the mapping has always modifier,
+ remove all the struct element mappings, which immediately
+ follow the GOMP_MAP_STRUCT map clause. */
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT)
+ {
+ HOST_WIDE_INT cnt = tree_to_shwi (OMP_CLAUSE_SIZE (c));
+ while (cnt--)
+ OMP_CLAUSE_CHAIN (c) = OMP_CLAUSE_CHAIN (OMP_CLAUSE_CHAIN (c));
+ }
+ }
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER
--- gcc/omp-low.c.jj 2015-08-24 14:32:06.000000000 +0200
+++ gcc/omp-low.c 2015-08-28 16:51:51.300696145 +0200
@@ -2074,6 +2074,12 @@ scan_sharing_clauses (tree clauses, omp_
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
{
+ if (TREE_CODE (decl) == COMPONENT_REF
+ || (TREE_CODE (decl) == INDIRECT_REF
+ && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+ == REFERENCE_TYPE)))
+ break;
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
@@ -13196,7 +13202,9 @@ lower_omp_target (gimple_stmt_iterator *
if (!DECL_P (var))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
- || !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
+ || (!OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
+ && (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_FIRSTPRIVATE_POINTER)))
map_cnt++;
continue;
}
@@ -13395,6 +13403,9 @@ lower_omp_target (gimple_stmt_iterator *
case OMP_CLAUSE_FROM:
nc = c;
ovar = OMP_CLAUSE_DECL (c);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ break;
if (!DECL_P (ovar))
{
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -13416,10 +13427,6 @@ lower_omp_target (gimple_stmt_iterator *
}
else
{
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c)
- == GOMP_MAP_FIRSTPRIVATE_POINTER)
- break;
if (DECL_SIZE (ovar)
&& TREE_CODE (DECL_SIZE (ovar)) != INTEGER_CST)
{
@@ -13880,10 +13887,19 @@ lower_omp_target (gimple_stmt_iterator *
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
{
location_t clause_loc = OMP_CLAUSE_LOCATION (c);
+ HOST_WIDE_INT offset = 0;
gcc_assert (prev);
var = OMP_CLAUSE_DECL (c);
- if (DECL_SIZE (var)
- && TREE_CODE (DECL_SIZE (var)) != INTEGER_CST)
+ if (TREE_CODE (var) == INDIRECT_REF
+ && TREE_CODE (TREE_OPERAND (var, 0)) == COMPONENT_REF)
+ var = TREE_OPERAND (var, 0);
+ if (TREE_CODE (var) == COMPONENT_REF)
+ {
+ var = get_addr_base_and_unit_offset (var, &offset);
+ gcc_assert (var != NULL_TREE && DECL_P (var));
+ }
+ else if (DECL_SIZE (var)
+ && TREE_CODE (DECL_SIZE (var)) != INTEGER_CST)
{
tree var2 = DECL_VALUE_EXPR (var);
gcc_assert (TREE_CODE (var2) == INDIRECT_REF);
@@ -13893,7 +13909,29 @@ lower_omp_target (gimple_stmt_iterator *
}
tree new_var = lookup_decl (var, ctx), x;
tree type = TREE_TYPE (new_var);
- bool is_ref = is_reference (var);
+ bool is_ref;
+ if (TREE_CODE (OMP_CLAUSE_DECL (c)) == INDIRECT_REF
+ && (TREE_CODE (TREE_OPERAND (OMP_CLAUSE_DECL (c), 0))
+ == COMPONENT_REF))
+ {
+ type = TREE_TYPE (TREE_OPERAND (OMP_CLAUSE_DECL (c), 0));
+ is_ref = true;
+ new_var = build2 (MEM_REF, type,
+ build_fold_addr_expr (new_var),
+ build_int_cst (build_pointer_type (type),
+ offset));
+ }
+ else if (TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF)
+ {
+ type = TREE_TYPE (OMP_CLAUSE_DECL (c));
+ is_ref = TREE_CODE (type) == REFERENCE_TYPE;
+ new_var = build2 (MEM_REF, type,
+ build_fold_addr_expr (new_var),
+ build_int_cst (build_pointer_type (type),
+ offset));
+ }
+ else
+ is_ref = is_reference (var);
bool ref_to_array = false;
if (is_ref)
{
--- gcc/c/c-typeck.c.jj 2015-07-31 16:58:09.000000000 +0200
+++ gcc/c/c-typeck.c 2015-08-27 18:53:04.122017251 +0200
@@ -11590,13 +11590,39 @@ c_finish_omp_cancellation_point (locatio
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
- bool &maybe_zero_len, unsigned int &first_non_one)
+ bool &maybe_zero_len, unsigned int &first_non_one,
+ bool is_omp)
{
tree ret, low_bound, length, type;
if (TREE_CODE (t) != TREE_LIST)
{
if (error_operand_p (t))
return error_mark_node;
+ ret = t;
+ if (TREE_CODE (t) == COMPONENT_REF
+ && is_omp
+ && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM))
+ {
+ 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)]);
+ return error_mark_node;
+ }
+ 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);
+ return error_mark_node;
+ }
+ t = TREE_OPERAND (t, 0);
+ }
+ }
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
if (DECL_P (t))
@@ -11617,11 +11643,11 @@ handle_omp_array_sections_1 (tree c, tre
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
return error_mark_node;
}
- return t;
+ return ret;
}
ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
- maybe_zero_len, first_non_one);
+ maybe_zero_len, first_non_one, is_omp);
if (ret == error_mark_node || ret == NULL_TREE)
return ret;
@@ -11856,7 +11882,8 @@ handle_omp_array_sections (tree c, bool
unsigned int first_non_one = 0;
auto_vec<tree, 10> types;
tree first = handle_omp_array_sections_1 (c, OMP_CLAUSE_DECL (c), types,
- maybe_zero_len, first_non_one);
+ maybe_zero_len, first_non_one,
+ is_omp);
if (first == error_mark_node)
return true;
if (first == NULL_TREE)
@@ -12027,7 +12054,9 @@ handle_omp_array_sections (tree c, bool
if (size)
size = c_fully_fold (size, false, NULL);
OMP_CLAUSE_SIZE (c) = size;
- if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || (TREE_CODE (t) == COMPONENT_REF
+ && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
return false;
gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DEVICEPTR);
if (is_omp)
@@ -12118,7 +12147,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, map_field_head;
+ bitmap_head aligned_head, map_head, map_field_head, generic_field_head;
tree c, t, type, *pc;
tree simdlen = NULL_TREE, safelen = NULL_TREE;
bool branch_seen = false;
@@ -12132,6 +12161,7 @@ c_finish_omp_clauses (tree clauses, bool
bitmap_initialize (&aligned_head, &bitmap_default_obstack);
bitmap_initialize (&map_head, &bitmap_default_obstack);
bitmap_initialize (&map_field_head, &bitmap_default_obstack);
+ bitmap_initialize (&generic_field_head, &bitmap_default_obstack);
for (pc = &clauses, c = clauses; c ; c = *pc)
{
@@ -12572,6 +12602,31 @@ c_finish_omp_clauses (tree clauses, bool
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
+ while (TREE_CODE (t) == ARRAY_REF)
+ t = TREE_OPERAND (t, 0);
+ if (TREE_CODE (t) == COMPONENT_REF
+ && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
+ {
+ while (TREE_CODE (t) == COMPONENT_REF)
+ t = TREE_OPERAND (t, 0);
+ if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ break;
+ if (bitmap_bit_p (&map_head, DECL_UID (t)))
+ {
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ error ("%qD appears more than once in motion"
+ " clauses", t);
+ else
+ error ("%qD appears more than once in map"
+ " clauses", t);
+ remove = true;
+ }
+ else
+ {
+ bitmap_set_bit (&map_head, DECL_UID (t));
+ bitmap_set_bit (&map_field_head, DECL_UID (t));
+ }
+ }
}
break;
}
@@ -12614,7 +12669,14 @@ c_finish_omp_clauses (tree clauses, bool
break;
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
{
- if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_POINTER))
+ {
+ if (bitmap_bit_p (&generic_field_head, DECL_UID (t)))
+ break;
+ }
+ else if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
break;
}
}
@@ -12648,6 +12710,23 @@ c_finish_omp_clauses (tree clauses, bool
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ {
+ if (bitmap_bit_p (&generic_head, DECL_UID (t))
+ || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+ {
+ error ("%qD appears more than once in data clauses", t);
+ remove = true;
+ }
+ else
+ {
+ bitmap_set_bit (&generic_head, DECL_UID (t));
+ if (t != OMP_CLAUSE_DECL (c)
+ && TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF)
+ bitmap_set_bit (&generic_field_head, DECL_UID (t));
+ }
+ }
else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
--- gcc/cp/parser.c.jj 2015-07-29 18:52:12.000000000 +0200
+++ gcc/cp/parser.c 2015-08-27 17:15:34.505155446 +0200
@@ -27950,10 +27950,22 @@ cp_parser_omp_var_list_no_open (cp_parse
decl = error_mark_node;
break;
}
- /* FALL THROUGH. */
+ /* FALLTHROUGH. */
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_FROM:
case OMP_CLAUSE_TO:
+ while (cp_lexer_next_token_is (parser->lexer, CPP_DOT))
+ {
+ location_t loc
+ = cp_lexer_peek_token (parser->lexer)->location;
+ cp_id_kind idk = CP_ID_KIND_NONE;
+ cp_lexer_consume_token (parser->lexer);
+ decl
+ = cp_parser_postfix_dot_deref_expression (parser, CPP_DOT,
+ decl, false,
+ &idk, loc);
+ }
+ /* FALLTHROUGH. */
case OMP_CLAUSE_DEPEND:
case OMP_CLAUSE_REDUCTION:
while (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_SQUARE))
--- gcc/cp/semantics.c.jj 2015-07-31 16:57:22.000000000 +0200
+++ gcc/cp/semantics.c 2015-08-28 13:59:51.033867196 +0200
@@ -4366,7 +4366,8 @@ omp_privatize_field (tree t)
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
- bool &maybe_zero_len, unsigned int &first_non_one)
+ bool &maybe_zero_len, unsigned int &first_non_one,
+ bool is_omp)
{
tree ret, low_bound, length, type;
if (TREE_CODE (t) != TREE_LIST)
@@ -4375,6 +4376,34 @@ handle_omp_array_sections_1 (tree c, tre
return error_mark_node;
if (type_dependent_expression_p (t))
return NULL_TREE;
+ if (REFERENCE_REF_P (t)
+ && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
+ t = TREE_OPERAND (t, 0);
+ ret = t;
+ if (TREE_CODE (t) == COMPONENT_REF
+ && is_omp
+ && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM))
+ {
+ 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)]);
+ return error_mark_node;
+ }
+ 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);
+ return error_mark_node;
+ }
+ t = TREE_OPERAND (t, 0);
+ }
+ }
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
if (processing_template_decl)
@@ -4406,15 +4435,15 @@ handle_omp_array_sections_1 (tree c, tre
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
return error_mark_node;
}
- t = convert_from_reference (t);
- return t;
+ ret = convert_from_reference (ret);
+ return ret;
}
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
&& TREE_CODE (TREE_CHAIN (t)) == FIELD_DECL)
TREE_CHAIN (t) = omp_privatize_field (TREE_CHAIN (t));
ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
- maybe_zero_len, first_non_one);
+ maybe_zero_len, first_non_one, is_omp);
if (ret == error_mark_node || ret == NULL_TREE)
return ret;
@@ -4656,7 +4685,8 @@ handle_omp_array_sections (tree c, bool
unsigned int first_non_one = 0;
auto_vec<tree, 10> types;
tree first = handle_omp_array_sections_1 (c, OMP_CLAUSE_DECL (c), types,
- maybe_zero_len, first_non_one);
+ maybe_zero_len, first_non_one,
+ is_omp);
if (first == error_mark_node)
return true;
if (first == NULL_TREE)
@@ -4824,7 +4854,9 @@ handle_omp_array_sections (tree c, bool
}
OMP_CLAUSE_DECL (c) = first;
OMP_CLAUSE_SIZE (c) = size;
- if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || (TREE_CODE (t) == COMPONENT_REF
+ && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
return false;
if (is_omp)
switch (OMP_CLAUSE_MAP_KIND (c))
@@ -5596,7 +5628,7 @@ tree
finish_omp_clauses (tree clauses, bool allow_fields, 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, generic_field_head;
tree c, t, *pc;
tree safelen = NULL_TREE;
bool branch_seen = false;
@@ -5608,6 +5640,8 @@ finish_omp_clauses (tree clauses, bool a
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);
+ bitmap_initialize (&generic_field_head, &bitmap_default_obstack);
for (pc = &clauses, c = clauses; c ; c = *pc)
{
@@ -6262,12 +6296,90 @@ finish_omp_clauses (tree clauses, bool a
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
+ while (TREE_CODE (t) == ARRAY_REF)
+ t = TREE_OPERAND (t, 0);
+ if (TREE_CODE (t) == COMPONENT_REF
+ && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
+ {
+ while (TREE_CODE (t) == COMPONENT_REF)
+ t = TREE_OPERAND (t, 0);
+ if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ break;
+ if (bitmap_bit_p (&map_head, DECL_UID (t)))
+ {
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ error ("%qD appears more than once in motion"
+ " clauses", t);
+ else
+ error ("%qD appears more than once in map"
+ " clauses", t);
+ remove = true;
+ }
+ else
+ {
+ bitmap_set_bit (&map_head, DECL_UID (t));
+ bitmap_set_bit (&map_field_head, DECL_UID (t));
+ }
+ }
}
break;
}
if (t == error_mark_node)
- remove = true;
- else if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
+ {
+ remove = true;
+ break;
+ }
+ if (REFERENCE_REF_P (t)
+ && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
+ t = TREE_OPERAND (t, 0);
+ if (TREE_CODE (t) == COMPONENT_REF
+ && allow_fields
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
+ {
+ if (type_dependent_expression_p (t))
+ break;
+ 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 (!cp_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 (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_POINTER))
+ {
+ if (bitmap_bit_p (&generic_field_head, DECL_UID (t)))
+ break;
+ }
+ else if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ break;
+ }
+ }
+ if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
if (processing_template_decl)
break;
@@ -6303,6 +6415,7 @@ finish_omp_clauses (tree clauses, bool a
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_FIRSTPRIVATE_POINTER)))
+ && t == OMP_CLAUSE_DECL (c)
&& !type_dependent_expression_p (t)
&& !cp_omp_mappable_type ((TREE_CODE (TREE_TYPE (t))
== REFERENCE_TYPE)
@@ -6314,6 +6427,27 @@ finish_omp_clauses (tree clauses, bool a
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ {
+ if (bitmap_bit_p (&generic_head, DECL_UID (t))
+ || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+ {
+ error ("%qD appears more than once in data clauses", t);
+ remove = true;
+ }
+ else
+ {
+ bitmap_set_bit (&generic_head, DECL_UID (t));
+ if (t != OMP_CLAUSE_DECL (c)
+ && (TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF
+ || (REFERENCE_REF_P (OMP_CLAUSE_DECL (c))
+ && (TREE_CODE (TREE_OPERAND (OMP_CLAUSE_DECL (c),
+ 0))
+ == COMPONENT_REF))))
+ bitmap_set_bit (&generic_field_head, DECL_UID (t));
+ }
+ }
else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
@@ -6323,7 +6457,12 @@ finish_omp_clauses (tree clauses, bool a
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/testsuite/c-c++-common/gomp/clauses-2.c.jj 2015-08-28 10:54:34.545144458 +0200
+++ gcc/testsuite/c-c++-common/gomp/clauses-2.c 2015-08-28 11:19:58.601066200 +0200
@@ -0,0 +1,53 @@
+struct S { int r; int *s; int t[10]; };
+void bar (int *);
+
+void
+foo (int *p, int q, struct S t, int i, int j, int k, int l)
+{
+ #pragma omp target map (q), firstprivate (q)
+ bar (&q);
+ #pragma omp target map (p[0]) firstprivate (p) /* { dg-error "appears more than once in data clauses" } */
+ bar (p);
+ #pragma omp target firstprivate (p), map (p[0]) /* { dg-error "appears more than once in data clauses" } */
+ bar (p);
+ #pragma omp target map (p[0]) map (p)
+ bar (p);
+ #pragma omp target map (p) , map (p[0])
+ bar (p);
+ #pragma omp target map (q) map (q) /* { dg-error "appears more than once in map clauses" } */
+ bar (&q);
+ #pragma omp target map (p[0]) map (p[0]) /* { dg-error "appears more than once in data clauses" } */
+ bar (p);
+ #pragma omp target map (t) map (t.r) /* { dg-error "appears more than once in map clauses" } */
+ bar (&t.r);
+ #pragma omp target map (t.r) map (t) /* { dg-error "appears more than once in map clauses" } */
+ bar (&t.r);
+ #pragma omp target map (t.r) map (t.r) /* { dg-error "appears more than once in map clauses" } */
+ bar (&t.r);
+ #pragma omp target firstprivate (t), map (t.r)
+ bar (&t.r);
+ #pragma omp target map (t.r) firstprivate (t)
+ bar (&t.r);
+ #pragma omp target map (t.s[0]) map (t)
+ bar (t.s);
+ #pragma omp target map (t) map(t.s[0])
+ bar (t.s);
+ #pragma omp target firstprivate (t) map (t.s[0]) /* { dg-error "appears more than once in data clauses" } */
+ bar (t.s);
+ #pragma omp target map (t.s[0]) firstprivate (t) /* { dg-error "appears more than once in data clauses" } */
+ bar (t.s);
+ #pragma omp target map (t.s[0]) map (t.s[2]) /* { dg-error "appears more than once in map clauses" } */
+ bar (t.s);
+ #pragma omp target map (t.t[0:2]) map (t.t[4:6]) /* { dg-error "appears more than once in map clauses" } */
+ bar (t.t);
+ #pragma omp target map (t.t[i:j]) map (t.t[k:l]) /* { dg-error "appears more than once in map clauses" } */
+ bar (t.t);
+ #pragma omp target map (t.s[0]) map (t.r)
+ bar (t.s);
+ #pragma omp target map (t.r) ,map (t.s[0])
+ bar (t.s);
+ #pragma omp target map (t.r) map (t) map (t.s[0]) firstprivate (t) /* { dg-error "appears more than once in map clauses" } */
+ bar (t.s); /* { dg-error "appears more than once in data clauses" "" { target *-*-* } 49 } */
+ #pragma omp target map (t) map (t.r) firstprivate (t) map (t.s[0]) /* { dg-error "appears more than once in map clauses" } */
+ bar (t.s); /* { dg-error "appears more than once in data clauses" "" { target *-*-* } 51 } */
+}
--- gcc/testsuite/c-c++-common/gomp/clauses-3.c.jj 2015-08-28 19:56:08.924530062 +0200
+++ gcc/testsuite/c-c++-common/gomp/clauses-3.c 2015-08-28 19:48:19.000000000 +0200
@@ -0,0 +1,23 @@
+struct T { int a; int *b; };
+struct S { int *s; char u; struct T v; long x; };
+
+void bar (int *);
+#pragma omp declare target to (bar)
+
+int
+main ()
+{
+ int a[10] = { 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 };
+ struct S s = { a, 5, { 6, a + 5 }, 99L };
+ #pragma omp target map (s.v.a, s.u, s.x)
+ ;
+ #pragma omp target map (s.v.a, s.u, s.x)
+ bar (&s.v.a);
+ #pragma omp target map (s.v.a) map (always, to: s.u) map (s.x)
+ ;
+ #pragma omp target map (s.s[0]) map (s.v.b[:3])
+ ;
+ #pragma omp target map (s.s[0]) map (s.v.b[:3])
+ bar (s.s);
+ return 0;
+}
--- libgomp/testsuite/libgomp.c++/target-10.C.jj 2015-08-28 10:57:13.898941691 +0200
+++ libgomp/testsuite/libgomp.c++/target-10.C 2015-08-28 10:57:03.822080985 +0200
@@ -0,0 +1 @@
+#include "../libgomp.c/target-21.c"
--- libgomp/testsuite/libgomp.c++/target-11.C.jj 2015-08-28 10:57:16.860900748 +0200
+++ libgomp/testsuite/libgomp.c++/target-11.C 2015-08-28 18:32:19.000000000 +0200
@@ -0,0 +1,62 @@
+extern "C" void abort ();
+struct T { int a; int *b; int c; char (&d)[10]; };
+struct S { int *s; char *u; struct T v; short *w; short *&x; };
+volatile int z;
+
+int
+main ()
+{
+ char d[10];
+ short *e;
+ int a[32], i;
+ char b[32];
+ short c[32];
+ for (i = 0; i < 32; i++)
+ {
+ a[i] = i;
+ b[i] = 32 + i;
+ c[i] = 64 + i;
+ }
+ for (i = 0; i < 10; i++)
+ d[i] = 17 + i;
+ e = c + 18;
+ struct S s = { a, b + 2, { 0, a + 16, 0, d }, c + 3, e };
+ int err = 0;
+ #pragma omp target map (to:s.v.b[0:z + 7], s.u[z + 1:z + 4]) \
+ map (tofrom:s.s[3:3], s.v.d[z + 1:z + 3]) \
+ map (from: s.w[z:4], s.x[1:3], err) private (i)
+ {
+ err = 0;
+ for (i = 0; i < 7; i++)
+ if (s.v.b[i] != 16 + i)
+ err = 1;
+ for (i = 1; i < 5; i++)
+ if (s.u[i] != 34 + i)
+ err = 1;
+ for (i = 3; i < 6; i++)
+ if (s.s[i] != i)
+ err = 1;
+ else
+ s.s[i] = 128 + i;
+ for (i = 1; i < 4; i++)
+ if (s.v.d[i] != 17 + i)
+ err = 1;
+ else
+ s.v.d[i] = 23 + i;
+ for (i = 0; i < 4; i++)
+ s.w[i] = 96 + i;
+ for (i = 1; i < 4; i++)
+ s.x[i] = 173 + i;
+ }
+ if (err)
+ abort ();
+ for (i = 0; i < 32; i++)
+ if (a[i] != ((i >= 3 && i < 6) ? 128 + i : i)
+ || b[i] != 32 + i
+ || c[i] != ((i >= 3 && i < 7) ? 93 + i : ((i >= 19 && i < 22) ? 155 + i : 64 + i)))
+ abort ();
+ for (i = 0; i < 10; i++)
+ if (d[i] != ((i >= 1 && i < 4) ? 23 + i : 17 + i))
+ abort ();
+ return 0;
+}
--- libgomp/testsuite/libgomp.c/target-21.c.jj 2015-07-31 17:32:56.000000000 +0200
+++ libgomp/testsuite/libgomp.c/target-21.c 2015-08-28 10:56:21.849661175 +0200
@@ -1,7 +1,12 @@
-extern void abort (void);
+extern
+#ifdef __cplusplus
+"C"
+#endif
+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; };
+struct S { int s; int u; struct T v; int x[10]; union U w; int y[10]; int z[10]; };
+volatile int z;
int
main ()
@@ -13,43 +18,66 @@ main ()
s.v.b.y = 3LL;
s.v.c = 19;
s.w.x = 4;
+ s.x[0] = 7;
+ s.x[1] = 8;
+ s.y[3] = 9;
+ s.y[4] = 10;
+ s.y[5] = 11;
int err = 0;
- #pragma omp target map (to:s.v.b, s.u) map (from: s.w, err)
+ #pragma omp target map (to:s.v.b, s.u, s.x[0:z + 2]) \
+ map (tofrom:s.y[3:3]) \
+ map (from: s.w, s.z[z + 1:z + 3], err)
{
err = 0;
- if (s.u != 1 || s.v.b.y != 3LL)
+ if (s.u != 1 || s.v.b.y != 3LL || s.x[0] != 7 || s.x[1] != 8
+ || s.y[3] != 9 || s.y[4] != 10 || s.y[5] != 11)
err = 1;
s.w.x = 6;
+ s.y[3] = 12;
+ s.y[4] = 13;
+ s.y[5] = 14;
+ s.z[1] = 15;
+ s.z[2] = 16;
+ s.z[3] = 17;
}
- if (err || s.w.x != 6)
+ if (err || s.w.x != 6 || s.y[3] != 12 || s.y[4] != 13 || s.y[5] != 14
+ || s.z[1] != 15 || s.z[2] != 16 || s.z[3] != 17)
abort ();
s.u++;
s.v.a++;
s.v.b.y++;
s.w.x++;
+ s.x[1] = 18;
+ s.z[0] = 19;
#pragma omp target data map (tofrom: s)
- #pragma omp target map (always to: s.w, err) map (alloc:s.u, s.v.b)
+ #pragma omp target map (always to: s.w, s.x[1], err) map (alloc:s.u, s.v.b, s.z[z:z + 1])
{
err = 0;
- if (s.u != 2 || s.v.b.y != 4LL || s.w.x != 7)
+ if (s.u != 2 || s.v.b.y != 4LL || s.w.x != 7 || s.x[1] != 18 || s.z[0] != 19)
err = 1;
s.w.x = 8;
+ s.x[1] = 20;
+ s.z[0] = 21;
}
- if (err || s.w.x != 8)
+ if (err || s.w.x != 8 || s.x[1] != 20 || s.z[0] != 21)
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)
+ s.x[0] = 22;
+ s.x[1] = 23;
+ #pragma omp target data map (from: s.w, s.x[0:2]) map (to: s.v.b, s.u)
+ #pragma omp target map (always to: s.w, s.x[0:2], err) map (alloc:s.u, s.v.b)
{
err = 0;
- if (s.u != 3 || s.v.b.y != 5LL || s.w.x != 9)
+ if (s.u != 3 || s.v.b.y != 5LL || s.w.x != 9 || s.x[0] != 22 || s.x[1] != 23)
err = 1;
s.w.x = 11;
+ s.x[0] = 24;
+ s.x[1] = 25;
}
- if (err || s.w.x != 11)
+ if (err || s.w.x != 11 || s.x[0] != 24 || s.x[1] != 25)
abort ();
return 0;
}
--- libgomp/testsuite/libgomp.c/target-22.c.jj 2015-08-27 13:13:09.999364928 +0200
+++ libgomp/testsuite/libgomp.c/target-22.c 2015-08-28 18:39:06.758874289 +0200
@@ -0,0 +1,51 @@
+extern void abort (void);
+struct T { int a; int *b; int c; };
+struct S { int *s; char *u; struct T v; short *w; };
+volatile int z;
+
+int
+main ()
+{
+ struct S s;
+ int a[32], i;
+ char b[32];
+ short c[32];
+ for (i = 0; i < 32; i++)
+ {
+ a[i] = i;
+ b[i] = 32 + i;
+ c[i] = 64 + i;
+ }
+ s.s = a;
+ s.u = b + 2;
+ s.v.b = a + 16;
+ s.w = c + 3;
+ int err = 0;
+ #pragma omp target map (to:s.v.b[0:z + 7], s.u[z + 1:z + 4]) \
+ map (tofrom:s.s[3:3]) \
+ map (from: s.w[z:4], err) private (i)
+ {
+ err = 0;
+ for (i = 0; i < 7; i++)
+ if (s.v.b[i] != 16 + i)
+ err = 1;
+ for (i = 1; i < 5; i++)
+ if (s.u[i] != 34 + i)
+ err = 1;
+ for (i = 3; i < 6; i++)
+ if (s.s[i] != i)
+ err = 1;
+ else
+ s.s[i] = 128 + i;
+ for (i = 0; i < 4; i++)
+ s.w[i] = 96 + i;
+ }
+ if (err)
+ abort ();
+ for (i = 0; i < 32; i++)
+ if (a[i] != ((i >= 3 && i < 6) ? 128 + i : i)
+ || b[i] != 32 + i
+ || c[i] != ((i >= 3 && i < 7) ? 93 + i : 64 + i))
+ abort ();
+ return 0;
+}
Jakub
^ permalink raw reply [flat|nested] 11+ messages in thread
* [gomp4.1] Structure element mapping support
2015-08-28 18:23 ` [gomp4.1] WIP: Structure " Jakub Jelinek
@ 2015-08-31 15:08 ` Jakub Jelinek
2015-09-02 11:21 ` Ilya Verbin
0 siblings, 1 reply; 11+ messages in thread
From: Jakub Jelinek @ 2015-08-31 15:08 UTC (permalink / raw)
To: Ilya Verbin; +Cc: gcc-patches
On Fri, Aug 28, 2015 at 08:13:35PM +0200, Jakub Jelinek wrote:
> Here is my current WIP on further structure element mapping support
> (so, structure element {pointer,reference to pointer,reference to array}
> based array sections, start of C++ support (still need to add tests for
> template instantiation and verify it works properly)).
> I have still pending questions on mapping of references (other than
> array sections) and structure element references pending, hope they will be
> responded to soon and will be able to commit this next week.
And here is the version I've committed. The C++ references (other than
array sections) aren't finished, as I haven't heard from omp-lang on this
topic yet.
Also, another known still broken case is zero length array section handling
on target enter data and target exit data constructs (apparently if only
zero length based array section appears in target enter data construct,
then we treat it as if that construct is exit data instead, plus
delete on zero length array sections is broken too).
For delete of zero length array sections we'll need a new map kind
in any case, for enter data vs. exit data distinction perhaps when we add
a flags parameter to hold e.g. the nowait flag, we can add the exit data
flag (vs. enter data) bit there too and stop using the heuristics.
And we are missing a testcase to test private/firstprivate clauses
on target construct with C++ data members (both normal and in template).
2015-08-31 Jakub Jelinek <jakub@redhat.com>
* gimplify.c (gimplify_scan_omp_clauses): Handle
struct element GOMP_MAP_FIRSTPRIVATE_POINTER.
(gimplify_adjust_omp_clauses): Add CODE argument.
Handle removal of GOMP_MAP_FIRSTPRIVATE_POINTER
struct elements for struct not seen in target body.
Handle removal of struct mapping if struct is not
seen in target body. Remove GOMP_MAP_STRUCT
map clause on OMP_TARGET_EXIT_DATA.
(gimplify_oacc_cache, gimplify_omp_parallel, gimplify_omp_task,
gimplify_omp_for, gimplify_omp_workshare, gimplify_omp_target_update,
gimplify_expr): Adjust callers.
* omp-low.c (scan_sharing_clauses): Handle struct element
GOMP_MAP_FIRSTPRIVATE_POINTER.
(lower_omp_target): Likewise.
gcc/c/
* c-typeck.c (handle_omp_array_sections_1): Add IS_OMP argument, pass
it down recursively. Handle struct element based array sections.
(handle_omp_array_sections): Adjust caller. Handle struct element
based array sections.
(c_finish_omp_clauses): Handle struct element based array sections.
Use generic_head instead of map_head for GOMP_MAP_FIRSTPRIVATE_POINTER
duplicate testing.
gcc/cp/
* parser.c (cp_parser_omp_var_list_no_open): Parse struct element
on map/to/from clauses.
(cp_parser_omp_clause_map): Fix up parsing of delete kind.
* pt.c (tsubst_expr): For OMP_TARGET{,_DATA} pass true instead of
false to allows_field.
* semantics.c (handle_omp_array_sections_1): Add IS_OMP argument,
pass it down recursively. Handle struct element based array sections.
(handle_omp_array_sections): Adjust caller. Handle struct element
based array sections.
(finish_omp_clauses): Handle struct element mappings and struct
element based array sections. Use generic_head instead of map_head
for GOMP_MAP_FIRSTPRIVATE_POINTER duplicate testing.
gcc/testsuite/
* c-c++-common/gomp/clauses-2.c: New test.
* c-c++-common/gomp/clauses-3.c: New test.
libgomp/
* target.c (GOMP_target_enter_exit_data): Allow GOMP_MAP_STRUCT
for enter data and handle it properly.
* testsuite/libgomp.c++/target-10.C: New test.
* testsuite/libgomp.c++/target-11.C: New test.
* testsuite/libgomp.c++/target-12.C: New test.
* testsuite/libgomp.c/target-21.c (z): New variable.
(struct S, main): Add tests for struct element array based array
sections.
* testsuite/libgomp.c/target-22.c: New test.
* testsuite/libgomp.c/target-23.c: New test.
--- gcc/gimplify.c.jj 2015-08-24 14:32:06.000000000 +0200
+++ gcc/gimplify.c 2015-08-31 14:52:32.804028967 +0200
@@ -6203,6 +6203,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;
+ tree *orig_list_p = list_p;
ctx = new_omp_context (region_type);
outer_ctx = ctx->outer_context;
@@ -6443,13 +6444,31 @@ gimplify_scan_omp_clauses (tree *list_p,
}
if (!DECL_P (decl))
{
+ tree d = decl, *pd;
+ if (TREE_CODE (d) == ARRAY_REF)
+ {
+ while (TREE_CODE (d) == ARRAY_REF)
+ d = TREE_OPERAND (d, 0);
+ if (TREE_CODE (d) == COMPONENT_REF
+ && TREE_CODE (TREE_TYPE (d)) == ARRAY_TYPE)
+ decl = d;
+ }
+ pd = &OMP_CLAUSE_DECL (c);
+ if (d == decl
+ && TREE_CODE (decl) == INDIRECT_REF
+ && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+ == REFERENCE_TYPE))
+ {
+ pd = &TREE_OPERAND (decl, 0);
+ decl = TREE_OPERAND (decl, 0);
+ }
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)
+ if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
== GS_ERROR)
{
remove = true;
@@ -6478,18 +6497,49 @@ gimplify_scan_omp_clauses (tree *list_p,
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);
+ tree base = OMP_CLAUSE_DECL (c);
+ while (TREE_CODE (base) == ARRAY_REF)
+ base = TREE_OPERAND (base, 0);
+ if (TREE_CODE (base) == INDIRECT_REF)
+ base = TREE_OPERAND (base, 0);
+ base = get_inner_reference (base, &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)
+ bool ptr = (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_POINTER);
+ if (n == NULL || (n->value & (ptr ? GOVD_PRIVATE
+ : GOVD_MAP)) == 0)
{
+ if (ptr)
+ {
+ tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_PRIVATE);
+ OMP_CLAUSE_DECL (c2) = decl;
+ OMP_CLAUSE_CHAIN (c2) = *orig_list_p;
+ *orig_list_p = c2;
+ if (struct_map_to_clause == NULL)
+ struct_map_to_clause = new hash_map<tree, tree>;
+ tree *osc;
+ if (n == NULL || (n->value & GOVD_MAP) == 0)
+ osc = NULL;
+ else
+ osc = struct_map_to_clause->get (decl);
+ if (osc == NULL)
+ struct_map_to_clause->put (decl,
+ tree_cons (NULL_TREE,
+ c,
+ NULL_TREE));
+ else
+ *osc = tree_cons (*osc, c, NULL_TREE);
+ flags = GOVD_PRIVATE | GOVD_EXPLICIT;
+ goto do_add_decl;
+ }
*list_p = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (*list_p, GOMP_MAP_STRUCT);
@@ -6508,6 +6558,9 @@ gimplify_scan_omp_clauses (tree *list_p,
else
{
tree *osc = struct_map_to_clause->get (decl), *sc;
+ tree *pt = NULL;
+ if (!ptr && TREE_CODE (*osc) == TREE_LIST)
+ osc = &TREE_PURPOSE (*osc);
if (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS)
n->value |= GOVD_SEEN;
offset_int o1, o2;
@@ -6517,25 +6570,58 @@ gimplify_scan_omp_clauses (tree *list_p,
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)
+ if (ptr)
+ pt = osc;
+ else
+ sc = &OMP_CLAUSE_CHAIN (*osc);
+ for (; ptr ? (*pt && (sc = &TREE_VALUE (*pt)))
+ : *sc != c;
+ ptr ? (pt = &TREE_CHAIN (*pt))
+ : (sc = &OMP_CLAUSE_CHAIN (*sc)))
+ if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) != COMPONENT_REF
+ && (TREE_CODE (OMP_CLAUSE_DECL (*sc))
+ != INDIRECT_REF)
+ && TREE_CODE (OMP_CLAUSE_DECL (*sc)) != ARRAY_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);
+ base = OMP_CLAUSE_DECL (*sc);
+ if (TREE_CODE (base) == ARRAY_REF)
+ {
+ while (TREE_CODE (base) == ARRAY_REF)
+ base = TREE_OPERAND (base, 0);
+ if (TREE_CODE (base) != COMPONENT_REF
+ || (TREE_CODE (TREE_TYPE (base))
+ != ARRAY_TYPE))
+ break;
+ }
+ else if (TREE_CODE (base) == INDIRECT_REF
+ && (TREE_CODE (TREE_OPERAND (base, 0))
+ == COMPONENT_REF)
+ && (TREE_CODE (TREE_TYPE
+ (TREE_OPERAND (base, 0)))
+ == REFERENCE_TYPE))
+ base = TREE_OPERAND (base, 0);
+ base = get_inner_reference (base, &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) == ARRAY_REF)
+ d1 = TREE_OPERAND (d1, 0);
+ while (TREE_CODE (d2) == ARRAY_REF)
+ d2 = TREE_OPERAND (d2, 0);
+ if (TREE_CODE (d1) == INDIRECT_REF)
+ d1 = TREE_OPERAND (d1, 0);
+ if (TREE_CODE (d2) == INDIRECT_REF)
+ d2 = TREE_OPERAND (d2, 0);
while (TREE_CODE (d1) == COMPONENT_REF)
if (TREE_CODE (d2) == COMPONENT_REF
&& TREE_OPERAND (d1, 1)
@@ -6564,6 +6650,12 @@ gimplify_scan_omp_clauses (tree *list_p,
|| (wi::eq_p (o1, o2) && bitpos < bitpos2))
break;
}
+ if (ptr)
+ {
+ if (!remove)
+ *pt = tree_cons (TREE_PURPOSE (*osc), c, *pt);
+ break;
+ }
if (!remove)
OMP_CLAUSE_SIZE (*osc)
= size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
@@ -7081,7 +7173,8 @@ gimplify_adjust_omp_clauses_1 (splay_tre
}
static void
-gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p)
+gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p,
+ enum tree_code code)
{
struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
tree c, decl;
@@ -7176,11 +7269,51 @@ gimplify_adjust_omp_clauses (gimple_seq
case OMP_CLAUSE_MAP:
decl = OMP_CLAUSE_DECL (c);
if (!DECL_P (decl))
- break;
+ {
+ if ((ctx->region_type & ORT_TARGET) != 0
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ {
+ if (TREE_CODE (decl) == INDIRECT_REF
+ && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+ == REFERENCE_TYPE))
+ decl = TREE_OPERAND (decl, 0);
+ if (TREE_CODE (decl) == COMPONENT_REF)
+ {
+ while (TREE_CODE (decl) == COMPONENT_REF)
+ decl = TREE_OPERAND (decl, 0);
+ if (DECL_P (decl))
+ {
+ n = splay_tree_lookup (ctx->variables,
+ (splay_tree_key) decl);
+ if (!(n->value & GOVD_SEEN))
+ remove = true;
+ }
+ }
+ }
+ break;
+ }
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
if ((ctx->region_type & ORT_TARGET) != 0
&& !(n->value & GOVD_SEEN)
- && !(OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS))
+ && ((OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS) == 0
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT))
+ {
+ remove = true;
+ /* For struct element mapping, if struct is never referenced
+ in target block and none of the mapping has always modifier,
+ remove all the struct element mappings, which immediately
+ follow the GOMP_MAP_STRUCT map clause. */
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT)
+ {
+ HOST_WIDE_INT cnt = tree_to_shwi (OMP_CLAUSE_SIZE (c));
+ while (cnt--)
+ OMP_CLAUSE_CHAIN (c)
+ = OMP_CLAUSE_CHAIN (OMP_CLAUSE_CHAIN (c));
+ }
+ }
+ else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
+ && code == OMP_TARGET_EXIT_DATA)
remove = true;
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
@@ -7337,7 +7470,7 @@ gimplify_oacc_cache (tree *expr_p, gimpl
gimplify_scan_omp_clauses (&OACC_CACHE_CLAUSES (expr), pre_p, ORT_WORKSHARE,
OACC_CACHE);
- gimplify_adjust_omp_clauses (pre_p, &OACC_CACHE_CLAUSES (expr));
+ gimplify_adjust_omp_clauses (pre_p, &OACC_CACHE_CLAUSES (expr), OACC_CACHE);
/* TODO: Do something sensible with this information. */
@@ -7369,7 +7502,8 @@ gimplify_omp_parallel (tree *expr_p, gim
else
pop_gimplify_context (NULL);
- gimplify_adjust_omp_clauses (pre_p, &OMP_PARALLEL_CLAUSES (expr));
+ gimplify_adjust_omp_clauses (pre_p, &OMP_PARALLEL_CLAUSES (expr),
+ OMP_PARALLEL);
g = gimple_build_omp_parallel (body,
OMP_PARALLEL_CLAUSES (expr),
@@ -7405,7 +7539,7 @@ gimplify_omp_task (tree *expr_p, gimple_
else
pop_gimplify_context (NULL);
- gimplify_adjust_omp_clauses (pre_p, &OMP_TASK_CLAUSES (expr));
+ gimplify_adjust_omp_clauses (pre_p, &OMP_TASK_CLAUSES (expr), OMP_TASK);
g = gimple_build_omp_task (body,
OMP_TASK_CLAUSES (expr),
@@ -7984,7 +8118,8 @@ gimplify_omp_for (tree *expr_p, gimple_s
TREE_OPERAND (TREE_OPERAND (t, 1), 0) = var;
}
- gimplify_adjust_omp_clauses (pre_p, &OMP_FOR_CLAUSES (orig_for_stmt));
+ gimplify_adjust_omp_clauses (pre_p, &OMP_FOR_CLAUSES (orig_for_stmt),
+ TREE_CODE (orig_for_stmt));
int kind;
switch (TREE_CODE (orig_for_stmt))
@@ -8236,7 +8371,7 @@ gimplify_omp_workshare (tree *expr_p, gi
}
else
gimplify_and_add (OMP_BODY (expr), &body);
- gimplify_adjust_omp_clauses (pre_p, &OMP_CLAUSES (expr));
+ gimplify_adjust_omp_clauses (pre_p, &OMP_CLAUSES (expr), TREE_CODE (expr));
switch (TREE_CODE (expr))
{
@@ -8312,7 +8447,8 @@ gimplify_omp_target_update (tree *expr_p
}
gimplify_scan_omp_clauses (&OMP_STANDALONE_CLAUSES (expr), pre_p,
ORT_WORKSHARE, TREE_CODE (expr));
- gimplify_adjust_omp_clauses (pre_p, &OMP_STANDALONE_CLAUSES (expr));
+ gimplify_adjust_omp_clauses (pre_p, &OMP_STANDALONE_CLAUSES (expr),
+ TREE_CODE (expr));
stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr));
gimplify_seq_add_stmt (pre_p, stmt);
@@ -9396,7 +9532,8 @@ gimplify_expr (tree *expr_p, gimple_seq
gimplify_scan_omp_clauses (&OMP_CRITICAL_CLAUSES (*expr_p),
pre_p, ORT_WORKSHARE, OMP_CRITICAL);
gimplify_adjust_omp_clauses (pre_p,
- &OMP_CRITICAL_CLAUSES (*expr_p));
+ &OMP_CRITICAL_CLAUSES (*expr_p),
+ OMP_CRITICAL);
g = gimple_build_omp_critical (body,
OMP_CRITICAL_NAME (*expr_p),
OMP_CRITICAL_CLAUSES (*expr_p));
--- gcc/omp-low.c.jj 2015-08-24 14:32:06.000000000 +0200
+++ gcc/omp-low.c 2015-08-28 16:51:51.300696145 +0200
@@ -2074,6 +2074,12 @@ scan_sharing_clauses (tree clauses, omp_
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
{
+ if (TREE_CODE (decl) == COMPONENT_REF
+ || (TREE_CODE (decl) == INDIRECT_REF
+ && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+ == REFERENCE_TYPE)))
+ break;
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
@@ -13196,7 +13202,9 @@ lower_omp_target (gimple_stmt_iterator *
if (!DECL_P (var))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
- || !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
+ || (!OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
+ && (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_FIRSTPRIVATE_POINTER)))
map_cnt++;
continue;
}
@@ -13395,6 +13403,9 @@ lower_omp_target (gimple_stmt_iterator *
case OMP_CLAUSE_FROM:
nc = c;
ovar = OMP_CLAUSE_DECL (c);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ break;
if (!DECL_P (ovar))
{
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -13416,10 +13427,6 @@ lower_omp_target (gimple_stmt_iterator *
}
else
{
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c)
- == GOMP_MAP_FIRSTPRIVATE_POINTER)
- break;
if (DECL_SIZE (ovar)
&& TREE_CODE (DECL_SIZE (ovar)) != INTEGER_CST)
{
@@ -13880,10 +13887,19 @@ lower_omp_target (gimple_stmt_iterator *
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
{
location_t clause_loc = OMP_CLAUSE_LOCATION (c);
+ HOST_WIDE_INT offset = 0;
gcc_assert (prev);
var = OMP_CLAUSE_DECL (c);
- if (DECL_SIZE (var)
- && TREE_CODE (DECL_SIZE (var)) != INTEGER_CST)
+ if (TREE_CODE (var) == INDIRECT_REF
+ && TREE_CODE (TREE_OPERAND (var, 0)) == COMPONENT_REF)
+ var = TREE_OPERAND (var, 0);
+ if (TREE_CODE (var) == COMPONENT_REF)
+ {
+ var = get_addr_base_and_unit_offset (var, &offset);
+ gcc_assert (var != NULL_TREE && DECL_P (var));
+ }
+ else if (DECL_SIZE (var)
+ && TREE_CODE (DECL_SIZE (var)) != INTEGER_CST)
{
tree var2 = DECL_VALUE_EXPR (var);
gcc_assert (TREE_CODE (var2) == INDIRECT_REF);
@@ -13893,7 +13909,29 @@ lower_omp_target (gimple_stmt_iterator *
}
tree new_var = lookup_decl (var, ctx), x;
tree type = TREE_TYPE (new_var);
- bool is_ref = is_reference (var);
+ bool is_ref;
+ if (TREE_CODE (OMP_CLAUSE_DECL (c)) == INDIRECT_REF
+ && (TREE_CODE (TREE_OPERAND (OMP_CLAUSE_DECL (c), 0))
+ == COMPONENT_REF))
+ {
+ type = TREE_TYPE (TREE_OPERAND (OMP_CLAUSE_DECL (c), 0));
+ is_ref = true;
+ new_var = build2 (MEM_REF, type,
+ build_fold_addr_expr (new_var),
+ build_int_cst (build_pointer_type (type),
+ offset));
+ }
+ else if (TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF)
+ {
+ type = TREE_TYPE (OMP_CLAUSE_DECL (c));
+ is_ref = TREE_CODE (type) == REFERENCE_TYPE;
+ new_var = build2 (MEM_REF, type,
+ build_fold_addr_expr (new_var),
+ build_int_cst (build_pointer_type (type),
+ offset));
+ }
+ else
+ is_ref = is_reference (var);
bool ref_to_array = false;
if (is_ref)
{
--- gcc/c/c-typeck.c.jj 2015-07-31 16:58:09.000000000 +0200
+++ gcc/c/c-typeck.c 2015-08-27 18:53:04.122017251 +0200
@@ -11590,13 +11590,39 @@ c_finish_omp_cancellation_point (locatio
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
- bool &maybe_zero_len, unsigned int &first_non_one)
+ bool &maybe_zero_len, unsigned int &first_non_one,
+ bool is_omp)
{
tree ret, low_bound, length, type;
if (TREE_CODE (t) != TREE_LIST)
{
if (error_operand_p (t))
return error_mark_node;
+ ret = t;
+ if (TREE_CODE (t) == COMPONENT_REF
+ && is_omp
+ && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM))
+ {
+ 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)]);
+ return error_mark_node;
+ }
+ 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);
+ return error_mark_node;
+ }
+ t = TREE_OPERAND (t, 0);
+ }
+ }
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
if (DECL_P (t))
@@ -11617,11 +11643,11 @@ handle_omp_array_sections_1 (tree c, tre
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
return error_mark_node;
}
- return t;
+ return ret;
}
ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
- maybe_zero_len, first_non_one);
+ maybe_zero_len, first_non_one, is_omp);
if (ret == error_mark_node || ret == NULL_TREE)
return ret;
@@ -11856,7 +11882,8 @@ handle_omp_array_sections (tree c, bool
unsigned int first_non_one = 0;
auto_vec<tree, 10> types;
tree first = handle_omp_array_sections_1 (c, OMP_CLAUSE_DECL (c), types,
- maybe_zero_len, first_non_one);
+ maybe_zero_len, first_non_one,
+ is_omp);
if (first == error_mark_node)
return true;
if (first == NULL_TREE)
@@ -12027,7 +12054,9 @@ handle_omp_array_sections (tree c, bool
if (size)
size = c_fully_fold (size, false, NULL);
OMP_CLAUSE_SIZE (c) = size;
- if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || (TREE_CODE (t) == COMPONENT_REF
+ && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
return false;
gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DEVICEPTR);
if (is_omp)
@@ -12118,7 +12147,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, map_field_head;
+ bitmap_head aligned_head, map_head, map_field_head, generic_field_head;
tree c, t, type, *pc;
tree simdlen = NULL_TREE, safelen = NULL_TREE;
bool branch_seen = false;
@@ -12132,6 +12161,7 @@ c_finish_omp_clauses (tree clauses, bool
bitmap_initialize (&aligned_head, &bitmap_default_obstack);
bitmap_initialize (&map_head, &bitmap_default_obstack);
bitmap_initialize (&map_field_head, &bitmap_default_obstack);
+ bitmap_initialize (&generic_field_head, &bitmap_default_obstack);
for (pc = &clauses, c = clauses; c ; c = *pc)
{
@@ -12572,6 +12602,31 @@ c_finish_omp_clauses (tree clauses, bool
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
+ while (TREE_CODE (t) == ARRAY_REF)
+ t = TREE_OPERAND (t, 0);
+ if (TREE_CODE (t) == COMPONENT_REF
+ && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
+ {
+ while (TREE_CODE (t) == COMPONENT_REF)
+ t = TREE_OPERAND (t, 0);
+ if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ break;
+ if (bitmap_bit_p (&map_head, DECL_UID (t)))
+ {
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ error ("%qD appears more than once in motion"
+ " clauses", t);
+ else
+ error ("%qD appears more than once in map"
+ " clauses", t);
+ remove = true;
+ }
+ else
+ {
+ bitmap_set_bit (&map_head, DECL_UID (t));
+ bitmap_set_bit (&map_field_head, DECL_UID (t));
+ }
+ }
}
break;
}
@@ -12614,7 +12669,14 @@ c_finish_omp_clauses (tree clauses, bool
break;
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
{
- if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_POINTER))
+ {
+ if (bitmap_bit_p (&generic_field_head, DECL_UID (t)))
+ break;
+ }
+ else if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
break;
}
}
@@ -12648,6 +12710,23 @@ c_finish_omp_clauses (tree clauses, bool
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ {
+ if (bitmap_bit_p (&generic_head, DECL_UID (t))
+ || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+ {
+ error ("%qD appears more than once in data clauses", t);
+ remove = true;
+ }
+ else
+ {
+ bitmap_set_bit (&generic_head, DECL_UID (t));
+ if (t != OMP_CLAUSE_DECL (c)
+ && TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF)
+ bitmap_set_bit (&generic_field_head, DECL_UID (t));
+ }
+ }
else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
--- gcc/cp/parser.c.jj 2015-07-29 18:52:12.000000000 +0200
+++ gcc/cp/parser.c 2015-08-31 16:04:58.607705130 +0200
@@ -27950,10 +27950,22 @@ cp_parser_omp_var_list_no_open (cp_parse
decl = error_mark_node;
break;
}
- /* FALL THROUGH. */
+ /* FALLTHROUGH. */
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_FROM:
case OMP_CLAUSE_TO:
+ while (cp_lexer_next_token_is (parser->lexer, CPP_DOT))
+ {
+ location_t loc
+ = cp_lexer_peek_token (parser->lexer)->location;
+ cp_id_kind idk = CP_ID_KIND_NONE;
+ cp_lexer_consume_token (parser->lexer);
+ decl
+ = cp_parser_postfix_dot_deref_expression (parser, CPP_DOT,
+ decl, false,
+ &idk, loc);
+ }
+ /* FALLTHROUGH. */
case OMP_CLAUSE_DEPEND:
case OMP_CLAUSE_REDUCTION:
while (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_SQUARE))
@@ -29655,7 +29667,9 @@ cp_parser_omp_clause_map (cp_parser *par
int nth = 2;
if (cp_lexer_peek_nth_token (parser->lexer, 2)->type == CPP_COMMA)
nth++;
- if (cp_lexer_peek_nth_token (parser->lexer, nth)->type == CPP_NAME
+ if ((cp_lexer_peek_nth_token (parser->lexer, nth)->type == CPP_NAME
+ || (cp_lexer_peek_nth_token (parser->lexer, nth)->keyword
+ == RID_DELETE))
&& (cp_lexer_peek_nth_token (parser->lexer, nth + 1)->type
== CPP_COLON))
{
@@ -29683,8 +29697,6 @@ cp_parser_omp_clause_map (cp_parser *par
kind = always ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM;
else if (strcmp ("release", p) == 0)
kind = GOMP_MAP_RELEASE;
- else if (strcmp ("delete", p) == 0)
- kind = GOMP_MAP_DELETE;
else
{
cp_parser_error (parser, "invalid map kind");
@@ -29696,6 +29708,13 @@ cp_parser_omp_clause_map (cp_parser *par
cp_lexer_consume_token (parser->lexer);
cp_lexer_consume_token (parser->lexer);
}
+ else if (cp_lexer_next_token_is_keyword (parser->lexer, RID_DELETE)
+ && cp_lexer_peek_nth_token (parser->lexer, 2)->type == CPP_COLON)
+ {
+ kind = GOMP_MAP_DELETE;
+ cp_lexer_consume_token (parser->lexer);
+ cp_lexer_consume_token (parser->lexer);
+ }
nlist = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list,
NULL);
--- gcc/cp/pt.c.jj 2015-07-16 17:56:41.000000000 +0200
+++ gcc/cp/pt.c 2015-08-31 11:48:54.628801176 +0200
@@ -14543,7 +14543,7 @@ tsubst_expr (tree t, tree args, tsubst_f
case OMP_TARGET_DATA:
case OMP_TARGET:
- tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, false,
+ tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, true,
args, complain, in_decl);
keep_next_level (true);
stmt = begin_omp_structured_block ();
@@ -14558,10 +14558,12 @@ tsubst_expr (tree t, tree args, tsubst_f
break;
case OMP_TARGET_UPDATE:
- tmp = tsubst_omp_clauses (OMP_TARGET_UPDATE_CLAUSES (t), false, false,
+ case OMP_TARGET_ENTER_DATA:
+ case OMP_TARGET_EXIT_DATA:
+ tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), false, true,
args, complain, in_decl);
t = copy_node (t);
- OMP_TARGET_UPDATE_CLAUSES (t) = tmp;
+ OMP_STANDALONE_CLAUSES (t) = tmp;
add_stmt (t);
break;
--- gcc/cp/semantics.c.jj 2015-07-31 16:57:22.000000000 +0200
+++ gcc/cp/semantics.c 2015-08-28 19:58:50.108378664 +0200
@@ -4366,7 +4366,8 @@ omp_privatize_field (tree t)
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
- bool &maybe_zero_len, unsigned int &first_non_one)
+ bool &maybe_zero_len, unsigned int &first_non_one,
+ bool is_omp)
{
tree ret, low_bound, length, type;
if (TREE_CODE (t) != TREE_LIST)
@@ -4375,6 +4376,34 @@ handle_omp_array_sections_1 (tree c, tre
return error_mark_node;
if (type_dependent_expression_p (t))
return NULL_TREE;
+ if (REFERENCE_REF_P (t)
+ && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
+ t = TREE_OPERAND (t, 0);
+ ret = t;
+ if (TREE_CODE (t) == COMPONENT_REF
+ && is_omp
+ && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM))
+ {
+ 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)]);
+ return error_mark_node;
+ }
+ 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);
+ return error_mark_node;
+ }
+ t = TREE_OPERAND (t, 0);
+ }
+ }
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
if (processing_template_decl)
@@ -4406,15 +4435,15 @@ handle_omp_array_sections_1 (tree c, tre
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
return error_mark_node;
}
- t = convert_from_reference (t);
- return t;
+ ret = convert_from_reference (ret);
+ return ret;
}
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
&& TREE_CODE (TREE_CHAIN (t)) == FIELD_DECL)
TREE_CHAIN (t) = omp_privatize_field (TREE_CHAIN (t));
ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
- maybe_zero_len, first_non_one);
+ maybe_zero_len, first_non_one, is_omp);
if (ret == error_mark_node || ret == NULL_TREE)
return ret;
@@ -4656,7 +4685,8 @@ handle_omp_array_sections (tree c, bool
unsigned int first_non_one = 0;
auto_vec<tree, 10> types;
tree first = handle_omp_array_sections_1 (c, OMP_CLAUSE_DECL (c), types,
- maybe_zero_len, first_non_one);
+ maybe_zero_len, first_non_one,
+ is_omp);
if (first == error_mark_node)
return true;
if (first == NULL_TREE)
@@ -4824,7 +4854,9 @@ handle_omp_array_sections (tree c, bool
}
OMP_CLAUSE_DECL (c) = first;
OMP_CLAUSE_SIZE (c) = size;
- if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || (TREE_CODE (t) == COMPONENT_REF
+ && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
return false;
if (is_omp)
switch (OMP_CLAUSE_MAP_KIND (c))
@@ -5596,7 +5628,7 @@ tree
finish_omp_clauses (tree clauses, bool allow_fields, 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, generic_field_head;
tree c, t, *pc;
tree safelen = NULL_TREE;
bool branch_seen = false;
@@ -5608,6 +5640,8 @@ finish_omp_clauses (tree clauses, bool a
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);
+ bitmap_initialize (&generic_field_head, &bitmap_default_obstack);
for (pc = &clauses, c = clauses; c ; c = *pc)
{
@@ -6262,12 +6296,90 @@ finish_omp_clauses (tree clauses, bool a
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
+ while (TREE_CODE (t) == ARRAY_REF)
+ t = TREE_OPERAND (t, 0);
+ if (TREE_CODE (t) == COMPONENT_REF
+ && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
+ {
+ while (TREE_CODE (t) == COMPONENT_REF)
+ t = TREE_OPERAND (t, 0);
+ if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ break;
+ if (bitmap_bit_p (&map_head, DECL_UID (t)))
+ {
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ error ("%qD appears more than once in motion"
+ " clauses", t);
+ else
+ error ("%qD appears more than once in map"
+ " clauses", t);
+ remove = true;
+ }
+ else
+ {
+ bitmap_set_bit (&map_head, DECL_UID (t));
+ bitmap_set_bit (&map_field_head, DECL_UID (t));
+ }
+ }
}
break;
}
if (t == error_mark_node)
- remove = true;
- else if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
+ {
+ remove = true;
+ break;
+ }
+ if (REFERENCE_REF_P (t)
+ && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
+ t = TREE_OPERAND (t, 0);
+ if (TREE_CODE (t) == COMPONENT_REF
+ && allow_fields
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
+ {
+ if (type_dependent_expression_p (t))
+ break;
+ 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 (!cp_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 (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_POINTER))
+ {
+ if (bitmap_bit_p (&generic_field_head, DECL_UID (t)))
+ break;
+ }
+ else if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ break;
+ }
+ }
+ if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
if (processing_template_decl)
break;
@@ -6303,6 +6415,7 @@ finish_omp_clauses (tree clauses, bool a
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_FIRSTPRIVATE_POINTER)))
+ && t == OMP_CLAUSE_DECL (c)
&& !type_dependent_expression_p (t)
&& !cp_omp_mappable_type ((TREE_CODE (TREE_TYPE (t))
== REFERENCE_TYPE)
@@ -6314,6 +6427,27 @@ finish_omp_clauses (tree clauses, bool a
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ {
+ if (bitmap_bit_p (&generic_head, DECL_UID (t))
+ || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+ {
+ error ("%qD appears more than once in data clauses", t);
+ remove = true;
+ }
+ else
+ {
+ bitmap_set_bit (&generic_head, DECL_UID (t));
+ if (t != OMP_CLAUSE_DECL (c)
+ && (TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF
+ || (REFERENCE_REF_P (OMP_CLAUSE_DECL (c))
+ && (TREE_CODE (TREE_OPERAND (OMP_CLAUSE_DECL (c),
+ 0))
+ == COMPONENT_REF))))
+ bitmap_set_bit (&generic_field_head, DECL_UID (t));
+ }
+ }
else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
@@ -6323,7 +6457,12 @@ finish_omp_clauses (tree clauses, bool a
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/testsuite/c-c++-common/gomp/clauses-2.c.jj 2015-08-28 10:54:34.545144458 +0200
+++ gcc/testsuite/c-c++-common/gomp/clauses-2.c 2015-08-28 11:19:58.601066200 +0200
@@ -0,0 +1,53 @@
+struct S { int r; int *s; int t[10]; };
+void bar (int *);
+
+void
+foo (int *p, int q, struct S t, int i, int j, int k, int l)
+{
+ #pragma omp target map (q), firstprivate (q)
+ bar (&q);
+ #pragma omp target map (p[0]) firstprivate (p) /* { dg-error "appears more than once in data clauses" } */
+ bar (p);
+ #pragma omp target firstprivate (p), map (p[0]) /* { dg-error "appears more than once in data clauses" } */
+ bar (p);
+ #pragma omp target map (p[0]) map (p)
+ bar (p);
+ #pragma omp target map (p) , map (p[0])
+ bar (p);
+ #pragma omp target map (q) map (q) /* { dg-error "appears more than once in map clauses" } */
+ bar (&q);
+ #pragma omp target map (p[0]) map (p[0]) /* { dg-error "appears more than once in data clauses" } */
+ bar (p);
+ #pragma omp target map (t) map (t.r) /* { dg-error "appears more than once in map clauses" } */
+ bar (&t.r);
+ #pragma omp target map (t.r) map (t) /* { dg-error "appears more than once in map clauses" } */
+ bar (&t.r);
+ #pragma omp target map (t.r) map (t.r) /* { dg-error "appears more than once in map clauses" } */
+ bar (&t.r);
+ #pragma omp target firstprivate (t), map (t.r)
+ bar (&t.r);
+ #pragma omp target map (t.r) firstprivate (t)
+ bar (&t.r);
+ #pragma omp target map (t.s[0]) map (t)
+ bar (t.s);
+ #pragma omp target map (t) map(t.s[0])
+ bar (t.s);
+ #pragma omp target firstprivate (t) map (t.s[0]) /* { dg-error "appears more than once in data clauses" } */
+ bar (t.s);
+ #pragma omp target map (t.s[0]) firstprivate (t) /* { dg-error "appears more than once in data clauses" } */
+ bar (t.s);
+ #pragma omp target map (t.s[0]) map (t.s[2]) /* { dg-error "appears more than once in map clauses" } */
+ bar (t.s);
+ #pragma omp target map (t.t[0:2]) map (t.t[4:6]) /* { dg-error "appears more than once in map clauses" } */
+ bar (t.t);
+ #pragma omp target map (t.t[i:j]) map (t.t[k:l]) /* { dg-error "appears more than once in map clauses" } */
+ bar (t.t);
+ #pragma omp target map (t.s[0]) map (t.r)
+ bar (t.s);
+ #pragma omp target map (t.r) ,map (t.s[0])
+ bar (t.s);
+ #pragma omp target map (t.r) map (t) map (t.s[0]) firstprivate (t) /* { dg-error "appears more than once in map clauses" } */
+ bar (t.s); /* { dg-error "appears more than once in data clauses" "" { target *-*-* } 49 } */
+ #pragma omp target map (t) map (t.r) firstprivate (t) map (t.s[0]) /* { dg-error "appears more than once in map clauses" } */
+ bar (t.s); /* { dg-error "appears more than once in data clauses" "" { target *-*-* } 51 } */
+}
--- gcc/testsuite/c-c++-common/gomp/clauses-3.c.jj 2015-08-28 19:56:08.924530062 +0200
+++ gcc/testsuite/c-c++-common/gomp/clauses-3.c 2015-08-28 19:48:19.000000000 +0200
@@ -0,0 +1,23 @@
+struct T { int a; int *b; };
+struct S { int *s; char u; struct T v; long x; };
+
+void bar (int *);
+#pragma omp declare target to (bar)
+
+int
+main ()
+{
+ int a[10] = { 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 };
+ struct S s = { a, 5, { 6, a + 5 }, 99L };
+ #pragma omp target map (s.v.a, s.u, s.x)
+ ;
+ #pragma omp target map (s.v.a, s.u, s.x)
+ bar (&s.v.a);
+ #pragma omp target map (s.v.a) map (always, to: s.u) map (s.x)
+ ;
+ #pragma omp target map (s.s[0]) map (s.v.b[:3])
+ ;
+ #pragma omp target map (s.s[0]) map (s.v.b[:3])
+ bar (s.s);
+ return 0;
+}
--- libgomp/target.c.jj 2015-07-31 16:55:38.000000000 +0200
+++ libgomp/target.c 2015-08-31 15:35:03.670073075 +0200
@@ -1465,7 +1465,8 @@ GOMP_target_enter_exit_data (int device,
if (kind == GOMP_MAP_ALLOC
|| kind == GOMP_MAP_TO
- || kind == GOMP_MAP_ALWAYS_TO)
+ || kind == GOMP_MAP_ALWAYS_TO
+ || kind == GOMP_MAP_STRUCT)
{
is_enter_data = true;
break;
@@ -1483,8 +1484,15 @@ GOMP_target_enter_exit_data (int device,
if (is_enter_data)
for (i = 0; i < mapnum; i++)
- gomp_map_vars (devicep, 1, &hostaddrs[i], NULL, &sizes[i], &kinds[i],
- true, GOMP_MAP_VARS_ENTER_DATA);
+ if ((kinds[i] & 0xff) == GOMP_MAP_STRUCT)
+ {
+ gomp_map_vars (devicep, sizes[i] + 1, &hostaddrs[i], NULL, &sizes[i],
+ &kinds[i], true, GOMP_MAP_VARS_ENTER_DATA);
+ i += sizes[i];
+ }
+ else
+ gomp_map_vars (devicep, 1, &hostaddrs[i], NULL, &sizes[i], &kinds[i],
+ true, GOMP_MAP_VARS_ENTER_DATA);
else
gomp_exit_data (devicep, mapnum, hostaddrs, sizes, kinds);
}
--- libgomp/testsuite/libgomp.c++/target-10.C.jj 2015-08-28 10:57:13.898941691 +0200
+++ libgomp/testsuite/libgomp.c++/target-10.C 2015-08-31 11:06:58.000000000 +0200
@@ -0,0 +1,154 @@
+extern "C" 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; T v; int x[10]; union U w; int y[10]; int z[10]; };
+volatile int z;
+
+template <typename R>
+void
+foo ()
+{
+ R s;
+ s.template s = 0;
+ s.u = 1;
+ s.v.a = 2;
+ s.v.b.y = 3LL;
+ s.v.c = 19;
+ s.w.x = 4;
+ s.template x[0] = 7;
+ s.x[1] = 8;
+ s.y[3] = 9;
+ s.y[4] = 10;
+ s.y[5] = 11;
+ int err = 0;
+ #pragma omp target map (to:s.template v.template b, s.u, s.x[0:z + 2]) \
+ map (tofrom:s.y[3:3]) \
+ map (from: s.w, s.template z[z + 1:z + 3], err)
+ {
+ err = 0;
+ if (s.u != 1 || s.v.b.y != 3LL || s.x[0] != 7 || s.x[1] != 8
+ || s.y[3] != 9 || s.y[4] != 10 || s.y[5] != 11)
+ err = 1;
+ s.w.x = 6;
+ s.y[3] = 12;
+ s.y[4] = 13;
+ s.y[5] = 14;
+ s.z[1] = 15;
+ s.z[2] = 16;
+ s.z[3] = 17;
+ }
+ if (err || s.w.x != 6 || s.y[3] != 12 || s.y[4] != 13 || s.y[5] != 14
+ || s.z[1] != 15 || s.z[2] != 16 || s.z[3] != 17)
+ abort ();
+ s.u++;
+ s.v.a++;
+ s.v.b.y++;
+ s.w.x++;
+ s.x[1] = 18;
+ s.z[0] = 19;
+ #pragma omp target data map (tofrom: s)
+ #pragma omp target map (always to: s.template w, s.x[1], err) map (alloc:s.u, s. template v.template b, s.z[z:z + 1])
+ {
+ err = 0;
+ if (s.u != 2 || s.v.b.y != 4LL || s.w.x != 7 || s.x[1] != 18 || s.z[0] != 19)
+ err = 1;
+ s.w.x = 8;
+ s.x[1] = 20;
+ s.z[0] = 21;
+ }
+ if (err || s.w.x != 8 || s.x[1] != 20 || s.z[0] != 21)
+ abort ();
+ s.u++;
+ s.v.a++;
+ s.v.b.y++;
+ s.w.x++;
+ s.x[0] = 22;
+ s.x[1] = 23;
+ #pragma omp target data map (from: s.w, s.x[0:2]) map (to: s.v.b, s.u)
+ #pragma omp target map (always to: s.w, s.x[0:2], err) map (alloc:s.u, s.v.b)
+ {
+ err = 0;
+ if (s.u != 3 || s.v.b.y != 5LL || s.w.x != 9 || s.x[0] != 22 || s.x[1] != 23)
+ err = 1;
+ s.w.x = 11;
+ s.x[0] = 24;
+ s.x[1] = 25;
+ }
+ if (err || s.w.x != 11 || s.x[0] != 24 || s.x[1] != 25)
+ abort ();
+}
+
+int
+main ()
+{
+ 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;
+ s.x[0] = 7;
+ s.x[1] = 8;
+ s.y[3] = 9;
+ s.y[4] = 10;
+ s.y[5] = 11;
+ int err = 0;
+ #pragma omp target map (to:s.v.b, s.u, s.x[0:z + 2]) \
+ map (tofrom:s.y[3:3]) \
+ map (from: s.w, s.z[z + 1:z + 3], err)
+ {
+ err = 0;
+ if (s.u != 1 || s.v.b.y != 3LL || s.x[0] != 7 || s.x[1] != 8
+ || s.y[3] != 9 || s.y[4] != 10 || s.y[5] != 11)
+ err = 1;
+ s.w.x = 6;
+ s.y[3] = 12;
+ s.y[4] = 13;
+ s.y[5] = 14;
+ s.z[1] = 15;
+ s.z[2] = 16;
+ s.z[3] = 17;
+ }
+ if (err || s.w.x != 6 || s.y[3] != 12 || s.y[4] != 13 || s.y[5] != 14
+ || s.z[1] != 15 || s.z[2] != 16 || s.z[3] != 17)
+ abort ();
+ s.u++;
+ s.v.a++;
+ s.v.b.y++;
+ s.w.x++;
+ s.x[1] = 18;
+ s.z[0] = 19;
+ #pragma omp target data map (tofrom: s)
+ #pragma omp target map (always to: s.w, s.x[1], err) map (alloc:s.u, s.v.b, s.z[z:z + 1])
+ {
+ err = 0;
+ if (s.u != 2 || s.v.b.y != 4LL || s.w.x != 7 || s.x[1] != 18 || s.z[0] != 19)
+ err = 1;
+ s.w.x = 8;
+ s.x[1] = 20;
+ s.z[0] = 21;
+ }
+ if (err || s.w.x != 8 || s.x[1] != 20 || s.z[0] != 21)
+ abort ();
+ s.u++;
+ s.v.a++;
+ s.v.b.y++;
+ s.w.x++;
+ s.x[0] = 22;
+ s.x[1] = 23;
+ #pragma omp target data map (from: s.w, s.x[0:2]) map (to: s.v.b, s.u)
+ #pragma omp target map (always to: s.w, s.x[0:2], err) map (alloc:s.u, s.v.b)
+ {
+ err = 0;
+ if (s.u != 3 || s.v.b.y != 5LL || s.w.x != 9 || s.x[0] != 22 || s.x[1] != 23)
+ err = 1;
+ s.w.x = 11;
+ s.x[0] = 24;
+ s.x[1] = 25;
+ }
+ if (err || s.w.x != 11 || s.x[0] != 24 || s.x[1] != 25)
+ abort ();
+ foo <S> ();
+ return 0;
+}
--- libgomp/testsuite/libgomp.c++/target-11.C.jj 2015-08-28 10:57:16.860900748 +0200
+++ libgomp/testsuite/libgomp.c++/target-11.C 2015-08-31 12:01:17.000000000 +0200
@@ -0,0 +1,121 @@
+extern "C" void abort ();
+struct T { int a; int *b; int c; char (&d)[10]; };
+struct S { int *s; char *u; T v; short *w; short *&x; };
+volatile int z;
+
+template <typename A, typename B, typename C, typename D>
+void
+foo ()
+{
+ A d[10];
+ B *e;
+ C a[32], i;
+ A b[32];
+ B c[32];
+ for (i = 0; i < 32; i++)
+ {
+ a[i] = i;
+ b[i] = 32 + i;
+ c[i] = 64 + i;
+ }
+ for (i = 0; i < 10; i++)
+ d[i] = 17 + i;
+ e = c + 18;
+ D s = { a, b + 2, { 0, a + 16, 0, d }, c + 3, e };
+ int err = 0;
+ #pragma omp target map (to:s.v.b[0:z + 7], s.template u[z + 1:z + 4]) \
+ map (tofrom:s.s[3:3], s. template v. template d[z + 1:z + 3]) \
+ map (from: s.w[z:4], s.x[1:3], err) private (i)
+ {
+ err = 0;
+ for (i = 0; i < 7; i++)
+ if (s.v.b[i] != 16 + i)
+ err = 1;
+ for (i = 1; i < 5; i++)
+ if (s.u[i] != 34 + i)
+ err = 1;
+ for (i = 3; i < 6; i++)
+ if (s.s[i] != i)
+ err = 1;
+ else
+ s.s[i] = 128 + i;
+ for (i = 1; i < 4; i++)
+ if (s.v.d[i] != 17 + i)
+ err = 1;
+ else
+ s.v.d[i] = 23 + i;
+ for (i = 0; i < 4; i++)
+ s.w[i] = 96 + i;
+ for (i = 1; i < 4; i++)
+ s.x[i] = 173 + i;
+ }
+ if (err)
+ abort ();
+ for (i = 0; i < 32; i++)
+ if (a[i] != ((i >= 3 && i < 6) ? 128 + i : i)
+ || b[i] != 32 + i
+ || c[i] != ((i >= 3 && i < 7) ? 93 + i : ((i >= 19 && i < 22) ? 155 + i : 64 + i)))
+ abort ();
+ for (i = 0; i < 10; i++)
+ if (d[i] != ((i >= 1 && i < 4) ? 23 + i : 17 + i))
+ abort ();
+}
+
+int
+main ()
+{
+ char d[10];
+ short *e;
+ int a[32], i;
+ char b[32];
+ short c[32];
+ for (i = 0; i < 32; i++)
+ {
+ a[i] = i;
+ b[i] = 32 + i;
+ c[i] = 64 + i;
+ }
+ for (i = 0; i < 10; i++)
+ d[i] = 17 + i;
+ e = c + 18;
+ S s = { a, b + 2, { 0, a + 16, 0, d }, c + 3, e };
+ int err = 0;
+ #pragma omp target map (to:s.v.b[0:z + 7], s.u[z + 1:z + 4]) \
+ map (tofrom:s.s[3:3], s.v.d[z + 1:z + 3]) \
+ map (from: s.w[z:4], s.x[1:3], err) private (i)
+ {
+ err = 0;
+ for (i = 0; i < 7; i++)
+ if (s.v.b[i] != 16 + i)
+ err = 1;
+ for (i = 1; i < 5; i++)
+ if (s.u[i] != 34 + i)
+ err = 1;
+ for (i = 3; i < 6; i++)
+ if (s.s[i] != i)
+ err = 1;
+ else
+ s.s[i] = 128 + i;
+ for (i = 1; i < 4; i++)
+ if (s.v.d[i] != 17 + i)
+ err = 1;
+ else
+ s.v.d[i] = 23 + i;
+ for (i = 0; i < 4; i++)
+ s.w[i] = 96 + i;
+ for (i = 1; i < 4; i++)
+ s.x[i] = 173 + i;
+ }
+ if (err)
+ abort ();
+ for (i = 0; i < 32; i++)
+ if (a[i] != ((i >= 3 && i < 6) ? 128 + i : i)
+ || b[i] != 32 + i
+ || c[i] != ((i >= 3 && i < 7) ? 93 + i : ((i >= 19 && i < 22) ? 155 + i : 64 + i)))
+ abort ();
+ for (i = 0; i < 10; i++)
+ if (d[i] != ((i >= 1 && i < 4) ? 23 + i : 17 + i))
+ abort ();
+ foo <char, short, int, S> ();
+ return 0;
+}
--- libgomp/testsuite/libgomp.c++/target-12.C.jj 2015-08-31 15:39:10.329714036 +0200
+++ libgomp/testsuite/libgomp.c++/target-12.C 2015-08-31 15:56:32.809545094 +0200
@@ -0,0 +1,93 @@
+extern "C" void abort (void);
+struct S { int s; int *u; int v[5]; };
+volatile int z;
+
+template <typename T>
+void
+foo ()
+{
+ int u[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }, err = 0;
+ T s = { 9, u + 3, { 10, 11, 12, 13, 14 } };
+ int *v = u + 4;
+ #pragma omp target enter data map (to: s.s, s.template u[0:5]) map (alloc: s.template v[1:3])
+ s.s++;
+ u[3]++;
+ s.v[1]++;
+ #pragma omp target update to (s.template s) to (s.u[0:2], s.v[1:3])
+ #pragma omp target map (alloc: s.s, s.v[1:3]) map (from: err)
+ {
+ err = 0;
+ if (s.s != 10 || s.v[1] != 12 || s.v[2] != 12 || s.v[3] != 13)
+ err = 1;
+ if (v[-1] != 4 || v[0] != 4 || v[1] != 5 || v[2] != 6 || v[3] != 7)
+ err = 1;
+ s.s++;
+ s.v[2] += 2;
+ v[-1] = 5;
+ v[3] = 9;
+ }
+ if (err)
+ abort ();
+ #pragma omp target map (alloc: s.u[0:5])
+ {
+ err = 0;
+ if (s.u[0] != 5 || s.u[1] != 4 || s.u[2] != 5 || s.u[3] != 6 || s.u[4] != 9)
+ err = 1;
+ s.u[1] = 12;
+ }
+ #pragma omp target update from (s.s, s.u[0:5]) from (s.v[1:3])
+ if (err || s.s != 11 || u[0] != 0 || u[1] != 1 || u[2] != 2 || u[3] != 5
+ || u[4] != 12 || u[5] != 5 || u[6] != 6 || u[7] != 9 || u[8] != 8
+ || u[9] != 9 || s.v[0] != 10 || s.v[1] != 12 || s.v[2] != 14
+ || s.v[3] != 13 || s.v[4] != 14)
+ abort ();
+ #pragma omp target exit data map (release: s.s)
+ #pragma omp target exit data map (release: s.u[0:5])
+ #pragma omp target exit data map (delete: s.v[1:3])
+ #pragma omp target exit data map (release: s.s)
+}
+
+int
+main ()
+{
+ int u[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }, err = 0;
+ S s = { 9, u + 3, { 10, 11, 12, 13, 14 } };
+ int *v = u + 4;
+ #pragma omp target enter data map (to: s.s, s.u[0:5]) map (alloc: s.v[1:3])
+ s.s++;
+ u[3]++;
+ s.v[1]++;
+ #pragma omp target update to (s.s) to (s.u[0:2], s.v[1:3])
+ #pragma omp target map (alloc: s.s, s.v[1:3]) map (from: err)
+ {
+ err = 0;
+ if (s.s != 10 || s.v[1] != 12 || s.v[2] != 12 || s.v[3] != 13)
+ err = 1;
+ if (v[-1] != 4 || v[0] != 4 || v[1] != 5 || v[2] != 6 || v[3] != 7)
+ err = 1;
+ s.s++;
+ s.v[2] += 2;
+ v[-1] = 5;
+ v[3] = 9;
+ }
+ if (err)
+ abort ();
+ #pragma omp target map (alloc: s.u[0:5])
+ {
+ err = 0;
+ if (s.u[0] != 5 || s.u[1] != 4 || s.u[2] != 5 || s.u[3] != 6 || s.u[4] != 9)
+ err = 1;
+ s.u[1] = 12;
+ }
+ #pragma omp target update from (s.s, s.u[0:5]) from (s.v[1:3])
+ if (err || s.s != 11 || u[0] != 0 || u[1] != 1 || u[2] != 2 || u[3] != 5
+ || u[4] != 12 || u[5] != 5 || u[6] != 6 || u[7] != 9 || u[8] != 8
+ || u[9] != 9 || s.v[0] != 10 || s.v[1] != 12 || s.v[2] != 14
+ || s.v[3] != 13 || s.v[4] != 14)
+ abort ();
+ #pragma omp target exit data map (release: s.s)
+ #pragma omp target exit data map (release: s.u[0:5])
+ #pragma omp target exit data map (always, delete: s.v[1:3])
+ #pragma omp target exit data map (release: s.s)
+ #pragma omp target exit data map (always delete : s.v[1:3])
+}
--- libgomp/testsuite/libgomp.c/target-21.c.jj 2015-07-31 17:32:56.000000000 +0200
+++ libgomp/testsuite/libgomp.c/target-21.c 2015-08-31 12:06:13.994068316 +0200
@@ -1,7 +1,8 @@
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; };
+struct S { int s; int u; struct T v; int x[10]; union U w; int y[10]; int z[10]; };
+volatile int z;
int
main ()
@@ -13,43 +14,66 @@ main ()
s.v.b.y = 3LL;
s.v.c = 19;
s.w.x = 4;
+ s.x[0] = 7;
+ s.x[1] = 8;
+ s.y[3] = 9;
+ s.y[4] = 10;
+ s.y[5] = 11;
int err = 0;
- #pragma omp target map (to:s.v.b, s.u) map (from: s.w, err)
+ #pragma omp target map (to:s.v.b, s.u, s.x[0:z + 2]) \
+ map (tofrom:s.y[3:3]) \
+ map (from: s.w, s.z[z + 1:z + 3], err)
{
err = 0;
- if (s.u != 1 || s.v.b.y != 3LL)
+ if (s.u != 1 || s.v.b.y != 3LL || s.x[0] != 7 || s.x[1] != 8
+ || s.y[3] != 9 || s.y[4] != 10 || s.y[5] != 11)
err = 1;
s.w.x = 6;
+ s.y[3] = 12;
+ s.y[4] = 13;
+ s.y[5] = 14;
+ s.z[1] = 15;
+ s.z[2] = 16;
+ s.z[3] = 17;
}
- if (err || s.w.x != 6)
+ if (err || s.w.x != 6 || s.y[3] != 12 || s.y[4] != 13 || s.y[5] != 14
+ || s.z[1] != 15 || s.z[2] != 16 || s.z[3] != 17)
abort ();
s.u++;
s.v.a++;
s.v.b.y++;
s.w.x++;
+ s.x[1] = 18;
+ s.z[0] = 19;
#pragma omp target data map (tofrom: s)
- #pragma omp target map (always to: s.w, err) map (alloc:s.u, s.v.b)
+ #pragma omp target map (always to: s.w, s.x[1], err) map (alloc:s.u, s.v.b, s.z[z:z + 1])
{
err = 0;
- if (s.u != 2 || s.v.b.y != 4LL || s.w.x != 7)
+ if (s.u != 2 || s.v.b.y != 4LL || s.w.x != 7 || s.x[1] != 18 || s.z[0] != 19)
err = 1;
s.w.x = 8;
+ s.x[1] = 20;
+ s.z[0] = 21;
}
- if (err || s.w.x != 8)
+ if (err || s.w.x != 8 || s.x[1] != 20 || s.z[0] != 21)
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)
+ s.x[0] = 22;
+ s.x[1] = 23;
+ #pragma omp target data map (from: s.w, s.x[0:2]) map (to: s.v.b, s.u)
+ #pragma omp target map (always to: s.w, s.x[0:2], err) map (alloc:s.u, s.v.b)
{
err = 0;
- if (s.u != 3 || s.v.b.y != 5LL || s.w.x != 9)
+ if (s.u != 3 || s.v.b.y != 5LL || s.w.x != 9 || s.x[0] != 22 || s.x[1] != 23)
err = 1;
s.w.x = 11;
+ s.x[0] = 24;
+ s.x[1] = 25;
}
- if (err || s.w.x != 11)
+ if (err || s.w.x != 11 || s.x[0] != 24 || s.x[1] != 25)
abort ();
return 0;
}
--- libgomp/testsuite/libgomp.c/target-22.c.jj 2015-08-27 13:13:09.999364928 +0200
+++ libgomp/testsuite/libgomp.c/target-22.c 2015-08-28 19:58:50.109378650 +0200
@@ -0,0 +1,51 @@
+extern void abort (void);
+struct T { int a; int *b; int c; };
+struct S { int *s; char *u; struct T v; short *w; };
+volatile int z;
+
+int
+main ()
+{
+ struct S s;
+ int a[32], i;
+ char b[32];
+ short c[32];
+ for (i = 0; i < 32; i++)
+ {
+ a[i] = i;
+ b[i] = 32 + i;
+ c[i] = 64 + i;
+ }
+ s.s = a;
+ s.u = b + 2;
+ s.v.b = a + 16;
+ s.w = c + 3;
+ int err = 0;
+ #pragma omp target map (to:s.v.b[0:z + 7], s.u[z + 1:z + 4]) \
+ map (tofrom:s.s[3:3]) \
+ map (from: s.w[z:4], err) private (i)
+ {
+ err = 0;
+ for (i = 0; i < 7; i++)
+ if (s.v.b[i] != 16 + i)
+ err = 1;
+ for (i = 1; i < 5; i++)
+ if (s.u[i] != 34 + i)
+ err = 1;
+ for (i = 3; i < 6; i++)
+ if (s.s[i] != i)
+ err = 1;
+ else
+ s.s[i] = 128 + i;
+ for (i = 0; i < 4; i++)
+ s.w[i] = 96 + i;
+ }
+ if (err)
+ abort ();
+ for (i = 0; i < 32; i++)
+ if (a[i] != ((i >= 3 && i < 6) ? 128 + i : i)
+ || b[i] != 32 + i
+ || c[i] != ((i >= 3 && i < 7) ? 93 + i : 64 + i))
+ abort ();
+ return 0;
+}
--- libgomp/testsuite/libgomp.c/target-23.c.jj 2015-08-31 14:09:40.386455884 +0200
+++ libgomp/testsuite/libgomp.c/target-23.c 2015-08-31 14:10:33.475729499 +0200
@@ -0,0 +1,48 @@
+extern void abort (void);
+struct S { int s; int *u; int v[5]; };
+volatile int z;
+
+int
+main ()
+{
+ int u[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }, err = 0;
+ struct S s = { 9, u + 3, { 10, 11, 12, 13, 14 } };
+ int *v = u + 4;
+ #pragma omp target enter data map (to: s.s, s.u[0:5]) map (alloc: s.v[1:3])
+ s.s++;
+ u[3]++;
+ s.v[1]++;
+ #pragma omp target update to (s.s) to (s.u[0:2], s.v[1:3])
+ #pragma omp target map (alloc: s.s, s.v[1:3]) map (from: err)
+ {
+ err = 0;
+ if (s.s != 10 || s.v[1] != 12 || s.v[2] != 12 || s.v[3] != 13)
+ err = 1;
+ if (v[-1] != 4 || v[0] != 4 || v[1] != 5 || v[2] != 6 || v[3] != 7)
+ err = 1;
+ s.s++;
+ s.v[2] += 2;
+ v[-1] = 5;
+ v[3] = 9;
+ }
+ if (err)
+ abort ();
+ #pragma omp target map (alloc: s.u[0:5])
+ {
+ err = 0;
+ if (s.u[0] != 5 || s.u[1] != 4 || s.u[2] != 5 || s.u[3] != 6 || s.u[4] != 9)
+ err = 1;
+ s.u[1] = 12;
+ }
+ #pragma omp target update from (s.s, s.u[0:5]) from (s.v[1:3])
+ if (err || s.s != 11 || u[0] != 0 || u[1] != 1 || u[2] != 2 || u[3] != 5
+ || u[4] != 12 || u[5] != 5 || u[6] != 6 || u[7] != 9 || u[8] != 8
+ || u[9] != 9 || s.v[0] != 10 || s.v[1] != 12 || s.v[2] != 14
+ || s.v[3] != 13 || s.v[4] != 14)
+ abort ();
+ #pragma omp target exit data map (release: s.s)
+ #pragma omp target exit data map (release: s.u[0:5])
+ #pragma omp target exit data map (delete: s.v[1:3])
+ #pragma omp target exit data map (release: s.s)
+ return 0;
+}
Jakub
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [gomp4.1] Structure element mapping support
2015-08-31 15:08 ` [gomp4.1] " Jakub Jelinek
@ 2015-09-02 11:21 ` Ilya Verbin
2015-09-02 15:59 ` [gomp4.1] Depend clause support for offloading Jakub Jelinek
0 siblings, 1 reply; 11+ messages in thread
From: Ilya Verbin @ 2015-09-02 11:21 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: gcc-patches
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;
^
-- Ilya
^ permalink raw reply [flat|nested] 11+ messages in thread
* [gomp4.1] Depend clause support for offloading
2015-09-02 11:21 ` Ilya Verbin
@ 2015-09-02 15:59 ` Jakub Jelinek
2015-09-03 14:19 ` Jakub Jelinek
0 siblings, 1 reply; 11+ messages in thread
From: Jakub Jelinek @ 2015-09-02 15:59 UTC (permalink / raw)
To: Ilya Verbin; +Cc: gcc-patches
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
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [gomp4.1] Depend clause support for offloading
2015-09-02 15:59 ` [gomp4.1] Depend clause support for offloading Jakub Jelinek
@ 2015-09-03 14:19 ` Jakub Jelinek
2015-09-03 17:41 ` Jakub Jelinek
2015-09-04 10:38 ` Jakub Jelinek
0 siblings, 2 replies; 11+ messages in thread
From: Jakub Jelinek @ 2015-09-03 14:19 UTC (permalink / raw)
To: Ilya Verbin; +Cc: gcc-patches
On Wed, Sep 02, 2015 at 05:58:54PM +0200, Jakub Jelinek wrote:
> 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).
Added testcase revealed two (small) issues, here is a fix for that together
with the testcase.
BTW, unless we want to add (at least now) support for running tasks in
between sending offloading target requests for memory allocation or data
movement and the offloading target signalizing their completion (supposedly
we'd better then be able to perform something like writev, merge as many
requests as possible into one metarequest and then await the completion of
it), I think at least for now we can ignore nowait on
target {update,{enter,exit} data} if depend clause is not also present
(on the library side).
I'll try to work on target {update,{enter,exit} data} nowait depend next
(in that case we need to copy the arrays and create some gomp_task).
2015-09-03 Jakub Jelinek <jakub@redhat.com>
* omp-low.c (lower_depend_clauses): Set TREE_ADDRESSABLE on array.
(lower_omp_target): Use gimple_omp_target_clauses_ptr instead of
gimple_omp_task_clauses_ptr.
* testsuite/libgomp.c/target-25.c: New test.
--- gcc/omp-low.c.jj 2015-09-02 15:13:13.000000000 +0200
+++ gcc/omp-low.c 2015-09-03 15:24:15.153716381 +0200
@@ -12975,6 +12975,7 @@ lower_depend_clauses (tree *pclauses, gi
}
tree type = build_array_type_nelts (ptr_type_node, n_in + n_out + 2);
tree array = create_tmp_var (type);
+ TREE_ADDRESSABLE (array) = 1;
tree r = build4 (ARRAY_REF, ptr_type_node, array, size_int (0), NULL_TREE,
NULL_TREE);
g = gimple_build_assign (r, build_int_cst (ptr_type_node, n_in + n_out));
@@ -13182,7 +13183,7 @@ lower_omp_target (gimple_stmt_iterator *
{
push_gimplify_context ();
dep_bind = gimple_build_bind (NULL, NULL, make_node (BLOCK));
- lower_depend_clauses (gimple_omp_task_clauses_ptr (stmt),
+ lower_depend_clauses (gimple_omp_target_clauses_ptr (stmt),
&dep_ilist, &dep_olist);
}
--- libgomp/testsuite/libgomp.c/target-25.c.jj 2015-09-03 15:02:34.130651945 +0200
+++ libgomp/testsuite/libgomp.c/target-25.c 2015-09-03 15:49:52.077362256 +0200
@@ -0,0 +1,84 @@
+#include <stdlib.h>
+#include <unistd.h>
+
+int
+main ()
+{
+ int x = 0, y = 0, z = 0, s = 11, t = 12, u = 13, w = 7, err;
+ #pragma omp parallel
+ #pragma omp single
+ {
+ #pragma omp task depend(in: x)
+ {
+ usleep (5000);
+ x = 1;
+ }
+ #pragma omp task depend(in: x)
+ {
+ usleep (6000);
+ y = 2;
+ }
+ #pragma omp task depend(out: z)
+ {
+ usleep (7000);
+ z = 3;
+ }
+ #pragma omp target map(tofrom: x) firstprivate (y) depend(inout: x, z)
+ err = (x != 1 || y != 2 || z != 3);
+ if (err)
+ abort ();
+ #pragma omp task depend(in: x)
+ {
+ usleep (5000);
+ x = 4;
+ }
+ #pragma omp task depend(in: x)
+ {
+ usleep (4000);
+ y = 5;
+ }
+ #pragma omp task depend(in: z)
+ {
+ usleep (3000);
+ z = 6;
+ }
+ #pragma omp target enter data nowait map (to: w)
+ #pragma omp target enter data depend (inout: x, z) map (to: x, y, z)
+ #pragma omp target map (alloc: x, y, z)
+ {
+ err = (x != 4 || y != 5 || z != 6);
+ x = 7;
+ y = 8;
+ z = 9;
+ }
+ if (err)
+ abort ();
+ #pragma omp taskwait
+ #pragma omp target map (alloc: w)
+ {
+ err = w != 7;
+ w = 17;
+ }
+ if (err)
+ abort ();
+ #pragma omp task depend(in: x)
+ {
+ usleep (2000);
+ s = 14;
+ }
+ #pragma omp task depend(in: x)
+ {
+ usleep (3000);
+ t = 15;
+ }
+ #pragma omp task depend(in: z)
+ {
+ usleep (4000);
+ u = 16;
+ }
+ #pragma omp target exit data depend (inout: x, z) map (from: x, y, z, w)
+ if (x != 7 || y != 8 || z != 9 || s != 14 || t != 15 || u != 16 || w != 17)
+ abort ();
+ }
+ return 0;
+}
Jakub
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [gomp4.1] Depend clause support for offloading
2015-09-03 14:19 ` Jakub Jelinek
@ 2015-09-03 17:41 ` Jakub Jelinek
2015-09-04 10:38 ` Jakub Jelinek
1 sibling, 0 replies; 11+ messages in thread
From: Jakub Jelinek @ 2015-09-03 17:41 UTC (permalink / raw)
To: Ilya Verbin, Aldy Hernandez; +Cc: gcc-patches
Hi!
FYI, I've merged trunk into the gomp-4_1-branch, it has been a while since
that has been done. make -C check RUNTESTFLAGS=gomp.exp and
make check-target-libgomp still pass without offloading and when offloading
to mic emul (the latter with the libgomp.c/for-5.c and libgomp.c++/for-13.C
LTO ICEs that have been failing for a while).
Jakub
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [gomp4.1] Depend clause support for offloading
2015-09-03 14:19 ` Jakub Jelinek
2015-09-03 17:41 ` Jakub Jelinek
@ 2015-09-04 10:38 ` Jakub Jelinek
1 sibling, 0 replies; 11+ messages in thread
From: Jakub Jelinek @ 2015-09-04 10:38 UTC (permalink / raw)
To: Ilya Verbin; +Cc: gcc-patches
On Thu, Sep 03, 2015 at 04:16:50PM +0200, Jakub Jelinek wrote:
> On Wed, Sep 02, 2015 at 05:58:54PM +0200, Jakub Jelinek wrote:
> > 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).
>
> Added testcase revealed two (small) issues, here is a fix for that together
> with the testcase.
There has been a bug in the testcase (missing map(from:err) in 3 places),
which hid a problem that on target constructs with depend clause (what about
just nowait?) we have to avoid using GOMP_FIRSTPRIVATE_INT or copy value
into temporary and take temporary's address for GOMP_FIRSTPRIVATE unless
we can prove other tasks can't modify the value while waiting for
dependencies (if it is addressable or shared with other threads/tasks,
then we have to use GOMP_FIRSTPRIVATE with address of the real variable,
so that if other tasks change it, we pick up the right values).
2015-09-04 Jakub Jelinek <jakub@redhat.com>
* omp-low.c (lower_omp_target): If target has depend
clauses, avoid using GOMP_MAP_FIRSTPRIVATE_INT unless the
var is non-addressable and private in the current task.
Even for GOMP_MAP_FIRSTPRIVATE, if the var is non-addressable,
but shared or threadprivate, take address of the shared var
rather than initializing a temporary with the current value.
* testsuite/libgomp.c/target-25.c (main): Add missing
map(from: err) clauses to target constructs.
--- gcc/omp-low.c.jj 2015-09-03 16:36:31.000000000 +0200
+++ gcc/omp-low.c 2015-09-04 11:34:45.512416693 +0200
@@ -13236,6 +13236,7 @@ lower_omp_target (gimple_stmt_iterator *
location_t loc = gimple_location (stmt);
bool offloaded, data_region;
unsigned int map_cnt = 0;
+ bool has_depend = false;
offloaded = is_gimple_omp_offloaded (stmt);
switch (gimple_omp_target_kind (stmt))
@@ -13268,6 +13269,7 @@ lower_omp_target (gimple_stmt_iterator *
dep_bind = gimple_build_bind (NULL, NULL, make_node (BLOCK));
lower_depend_clauses (gimple_omp_target_clauses_ptr (stmt),
&dep_ilist, &dep_olist);
+ has_depend = true;
}
tgt_bind = NULL;
@@ -13719,9 +13721,44 @@ lower_omp_target (gimple_stmt_iterator *
type = TREE_TYPE (ovar);
if (is_reference (ovar))
type = TREE_TYPE (type);
+ bool use_firstprivate_int, force_addr;
+ use_firstprivate_int = false;
+ force_addr = false;
if ((INTEGRAL_TYPE_P (type)
- && TYPE_PRECISION (type) <= POINTER_SIZE)
+ && TYPE_PRECISION (type) <= POINTER_SIZE)
|| TREE_CODE (type) == POINTER_TYPE)
+ use_firstprivate_int = true;
+ if (has_depend)
+ {
+ if (is_reference (var))
+ use_firstprivate_int = false;
+ else if (is_gimple_reg (var))
+ {
+ if (DECL_HAS_VALUE_EXPR_P (var))
+ {
+ tree v = get_base_address (var);
+ if (DECL_P (v) && TREE_ADDRESSABLE (v))
+ {
+ use_firstprivate_int = false;
+ force_addr = true;
+ }
+ else
+ switch (TREE_CODE (v))
+ {
+ case INDIRECT_REF:
+ case MEM_REF:
+ use_firstprivate_int = false;
+ force_addr = true;
+ break;
+ default:
+ break;
+ }
+ }
+ }
+ else
+ use_firstprivate_int = false;
+ }
+ if (use_firstprivate_int)
{
tkind = GOMP_MAP_FIRSTPRIVATE_INT;
tree t = var;
@@ -13734,7 +13771,7 @@ lower_omp_target (gimple_stmt_iterator *
}
else if (is_reference (var))
gimplify_assign (x, var, &ilist);
- else if (is_gimple_reg (var))
+ else if (!force_addr && is_gimple_reg (var))
{
tree avar = create_tmp_var (TREE_TYPE (var));
mark_addressable (avar);
@@ -13867,9 +13904,40 @@ lower_omp_target (gimple_stmt_iterator *
type = TREE_TYPE (var);
if (is_reference (var))
type = TREE_TYPE (type);
+ bool use_firstprivate_int;
+ use_firstprivate_int = false;
if ((INTEGRAL_TYPE_P (type)
&& TYPE_PRECISION (type) <= POINTER_SIZE)
|| TREE_CODE (type) == POINTER_TYPE)
+ use_firstprivate_int = true;
+ if (has_depend)
+ {
+ tree v = lookup_decl_in_outer_ctx (var, ctx);
+ if (is_reference (v))
+ use_firstprivate_int = false;
+ else if (is_gimple_reg (v))
+ {
+ if (DECL_HAS_VALUE_EXPR_P (v))
+ {
+ v = get_base_address (v);
+ if (DECL_P (v) && TREE_ADDRESSABLE (v))
+ use_firstprivate_int = false;
+ else
+ switch (TREE_CODE (v))
+ {
+ case INDIRECT_REF:
+ case MEM_REF:
+ use_firstprivate_int = false;
+ break;
+ default:
+ break;
+ }
+ }
+ }
+ else
+ use_firstprivate_int = false;
+ }
+ if (use_firstprivate_int)
{
x = build_receiver_ref (var, false, ctx);
if (TREE_CODE (type) != POINTER_TYPE)
--- libgomp/testsuite/libgomp.c/target-25.c.jj 2015-09-04 10:41:52.371881507 +0200
+++ libgomp/testsuite/libgomp.c/target-25.c 2015-09-04 10:39:16.000000000 +0200
@@ -23,7 +23,7 @@ main ()
usleep (7000);
z = 3;
}
- #pragma omp target map(tofrom: x) firstprivate (y) depend(inout: x, z)
+ #pragma omp target map(tofrom: x) map(from: err) firstprivate (y) depend(inout: x, z)
err = (x != 1 || y != 2 || z != 3);
if (err)
abort ();
@@ -44,7 +44,7 @@ main ()
}
#pragma omp target enter data nowait map (to: w)
#pragma omp target enter data depend (inout: x, z) map (to: x, y, z)
- #pragma omp target map (alloc: x, y, z)
+ #pragma omp target map (alloc: x, y, z) map(from: err)
{
err = (x != 4 || y != 5 || z != 6);
x = 7;
@@ -54,7 +54,7 @@ main ()
if (err)
abort ();
#pragma omp taskwait
- #pragma omp target map (alloc: w)
+ #pragma omp target map (alloc: w) map(from: err)
{
err = w != 7;
w = 17;
Jakub
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [gomp4.1] Start of structure element mapping support
2015-07-31 16:28 [gomp4.1] Start of structure element mapping support Jakub Jelinek
2015-08-28 18:23 ` [gomp4.1] WIP: Structure " Jakub Jelinek
@ 2019-10-16 13:35 ` Thomas Schwinge
2019-10-16 17:46 ` Jakub Jelinek
1 sibling, 1 reply; 11+ messages in thread
From: Thomas Schwinge @ 2019-10-16 13:35 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: gcc-patches, Julian Brown
[-- Attachment #1: Type: text/plain, Size: 4159 bytes --]
Hi Jakub!
Stumbled over this while reviewing Julian's "Factor out duplicate code in
gimplify_scan_omp_clauses":
On 2015-07-31T18:16:10+0200, Jakub Jelinek <jakub@redhat.com> wrote:
> This patch is the start of implementation of struct element mapping.
Not quite the same, but similar code is still present in GCC trunk.
> --- gcc/gimplify.c.jj 2015-07-31 16:55:01.482411392 +0200
> +++ gcc/gimplify.c 2015-07-31 16:57:22.307320290 +0200
> + tree offset;
Here we define '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);
..., which here gets writte to...
> + gcc_assert (base == decl
> + && (offset == NULL_TREE
> + || TREE_CODE (offset) == INTEGER_CST));
..., and here gets checked...
> +
> + splay_tree_node n
> + = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
> + if (n == NULL || (n->value & GOVD_MAP) == 0)
> + {
> + [...]
> + }
> + 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);
..., and here used.
> + 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;
Here we define 'offset2'...
> + HOST_WIDE_INT bitsize2, bitpos2;
> + base = get_inner_reference (OMP_CLAUSE_DECL (*sc),
> + &bitsize2, &bitpos2,
> + &offset2, &mode,
> + &unsignedp, &volatilep,
> + false);
..., which here gets writte to...
> + if (base != decl)
> + break;
> + gcc_assert (offset == NULL_TREE
> + || TREE_CODE (offset) == INTEGER_CST);
..., but here we again check 'offset', not 'offset2'...
> + 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);
.., but here again we use '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;
Should the second highlighted 'gcc_assert' be changed as follows,
suitably adapted for current GCC trunk, of course? (Not yet tested.) If
approving such a patch, please respond with "Reviewed-by: NAME <EMAIL>"
so that your effort will be recorded in the commit log, see
<https://gcc.gnu.org/wiki/Reviewed-by>.
- gcc_assert (offset == NULL_TREE
- || TREE_CODE (offset) == INTEGER_CST);
+ gcc_assert (offset2 == NULL_TREE
+ || TREE_CODE (offset2) == INTEGER_CST);
Grüße
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 658 bytes --]
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [gomp4.1] Start of structure element mapping support
2019-10-16 13:35 ` [gomp4.1] Start of structure element mapping support Thomas Schwinge
@ 2019-10-16 17:46 ` Jakub Jelinek
2019-11-11 9:13 ` Thomas Schwinge
0 siblings, 1 reply; 11+ messages in thread
From: Jakub Jelinek @ 2019-10-16 17:46 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: gcc-patches, Julian Brown
On Wed, Oct 16, 2019 at 03:22:52PM +0200, Thomas Schwinge wrote:
> Stumbled over this while reviewing Julian's "Factor out duplicate code in
> gimplify_scan_omp_clauses":
> ..., which here gets writte to...
>
> > + if (base != decl)
> > + break;
> > + gcc_assert (offset == NULL_TREE
> > + || TREE_CODE (offset) == INTEGER_CST);
>
> ..., but here we again check 'offset', not 'offset2'...
Yes, it indeed should be offset2 == NULL_TREE and
TREE_CODE (offset2) == INTEGER_CST, thanks for catching that.
> Should the second highlighted 'gcc_assert' be changed as follows,
> suitably adapted for current GCC trunk, of course? (Not yet tested.) If
> approving such a patch, please respond with "Reviewed-by: NAME <EMAIL>"
> so that your effort will be recorded in the commit log, see
> <https://gcc.gnu.org/wiki/Reviewed-by>.
>
> - gcc_assert (offset == NULL_TREE
> - || TREE_CODE (offset) == INTEGER_CST);
> + gcc_assert (offset2 == NULL_TREE
> + || TREE_CODE (offset2) == INTEGER_CST);
Preapproved for trunk if it passes bootstrap/regtest.
Jakub
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [gomp4.1] Start of structure element mapping support
2019-10-16 17:46 ` Jakub Jelinek
@ 2019-11-11 9:13 ` Thomas Schwinge
0 siblings, 0 replies; 11+ messages in thread
From: Thomas Schwinge @ 2019-11-11 9:13 UTC (permalink / raw)
To: Jakub Jelinek, gcc-patches; +Cc: Julian Brown
[-- Attachment #1.1: Type: text/plain, Size: 896 bytes --]
Hi!
On 2019-10-16T18:52:55+0200, Jakub Jelinek <jakub@redhat.com> wrote:
> On Wed, Oct 16, 2019 at 03:22:52PM +0200, Thomas Schwinge wrote:
>> Stumbled over this while reviewing Julian's "Factor out duplicate code in
>> gimplify_scan_omp_clauses":
>
>> ..., which here gets writte to...
>>
>> > + if (base != decl)
>> > + break;
>> > + gcc_assert (offset == NULL_TREE
>> > + || TREE_CODE (offset) == INTEGER_CST);
>>
>> ..., but here we again check 'offset', not 'offset2'...
>
> Yes, it indeed [...]
Thanks. See attached; committed "Assert 'offset2' instead of 'offset' in
'gcc/gimplify.c:gimplify_scan_omp_clauses'" to trunk in r278038,
gcc-9-branch in r278039, gcc-8-branch in r278040, gcc-7-branch (slightly
different patch) omitted as that one's frozen for the final release, and
this fix isn't important enough.
Grüße
Thomas
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #1.2: 0001-Assert-offset2-instead-of-offset-in-gcc-gimpli.trunk.patch --]
[-- Type: text/x-diff, Size: 1637 bytes --]
From c88bb56e02340083d10a728c2ca05748df5c6218 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Mon, 11 Nov 2019 08:18:46 +0000
Subject: [PATCH] Assert 'offset2' instead of 'offset' in
'gcc/gimplify.c:gimplify_scan_omp_clauses'
... to fix a long-time typo/copy'n'past-o.
gcc/
* gimplify.c (gimplify_scan_omp_clauses): Assert 'offset2' instead
of 'offset'.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@278038 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog | 3 +++
gcc/gimplify.c | 4 ++--
2 files changed, 5 insertions(+), 2 deletions(-)
diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index a1e928bf804d..ca0cebc20180 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,5 +1,8 @@
2019-11-11 Thomas Schwinge <thomas@codesourcery.com>
+ * gimplify.c (gimplify_scan_omp_clauses): Assert 'offset2' instead
+ of 'offset'.
+
* Makefile.in (LANG_CONFIGUREFRAGS): Define.
(config.status): Use/depend on it.
* configure.ac (all_lang_configurefrags): Track, 'AC_SUBST'.
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 5fa0ba6dda60..2bc41cf98ae9 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8890,8 +8890,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
break;
if (scp)
continue;
- gcc_assert (offset == NULL_TREE
- || poly_int_tree_p (offset));
+ gcc_assert (offset2 == NULL_TREE
+ || poly_int_tree_p (offset2));
tree d1 = OMP_CLAUSE_DECL (*sc);
tree d2 = OMP_CLAUSE_DECL (c);
while (TREE_CODE (d1) == ARRAY_REF)
--
2.17.1
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #1.3: 0001-Assert-offset2-instead-of-offset-in-gcc.gcc-9-branch.patch --]
[-- Type: text/x-diff, Size: 1670 bytes --]
From a3174355911ea11b2e1f2786a7bffe2c5489e128 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Mon, 11 Nov 2019 08:19:23 +0000
Subject: [PATCH] Assert 'offset2' instead of 'offset' in
'gcc/gimplify.c:gimplify_scan_omp_clauses'
... to fix a long-time typo/copy'n'past-o.
gcc/
* gimplify.c (gimplify_scan_omp_clauses): Assert 'offset2' instead
of 'offset'.
Backport from trunk r278038.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gcc-9-branch@278039 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog | 4 ++++
gcc/gimplify.c | 4 ++--
2 files changed, 6 insertions(+), 2 deletions(-)
diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 52d867db50a1..fc19dae278d4 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,5 +1,9 @@
2019-11-11 Thomas Schwinge <thomas@codesourcery.com>
+ Backport from trunk:
+ * gimplify.c (gimplify_scan_omp_clauses): Assert 'offset2' instead
+ of 'offset'.
+
Backport from trunk:
* Makefile.in (LANG_CONFIGUREFRAGS): Define.
(config.status): Use/depend on it.
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 28e13a710215..a0cb6c402bcf 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8718,8 +8718,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
break;
if (scp)
continue;
- gcc_assert (offset == NULL_TREE
- || poly_int_tree_p (offset));
+ gcc_assert (offset2 == NULL_TREE
+ || poly_int_tree_p (offset2));
tree d1 = OMP_CLAUSE_DECL (*sc);
tree d2 = OMP_CLAUSE_DECL (c);
while (TREE_CODE (d1) == ARRAY_REF)
--
2.17.1
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #1.4: 0001-Assert-offset2-instead-of-offset-in-gcc.gcc-8-branch.patch --]
[-- Type: text/x-diff, Size: 1670 bytes --]
From 2c1fdb6ef7fc5fbde790143c9c58b0f10241f9d8 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Mon, 11 Nov 2019 08:19:41 +0000
Subject: [PATCH] Assert 'offset2' instead of 'offset' in
'gcc/gimplify.c:gimplify_scan_omp_clauses'
... to fix a long-time typo/copy'n'past-o.
gcc/
* gimplify.c (gimplify_scan_omp_clauses): Assert 'offset2' instead
of 'offset'.
Backport from trunk r278038.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gcc-8-branch@278040 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog | 4 ++++
gcc/gimplify.c | 4 ++--
2 files changed, 6 insertions(+), 2 deletions(-)
diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 80bf2e121c63..29441ada5dd3 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,5 +1,9 @@
2019-11-11 Thomas Schwinge <thomas@codesourcery.com>
+ Backport from trunk:
+ * gimplify.c (gimplify_scan_omp_clauses): Assert 'offset2' instead
+ of 'offset'.
+
Backport from trunk:
* Makefile.in (LANG_CONFIGUREFRAGS): Define.
(config.status): Use/depend on it.
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index fcb50c8e0476..3bfd0692103b 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8181,8 +8181,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
break;
if (scp)
continue;
- gcc_assert (offset == NULL_TREE
- || poly_int_tree_p (offset));
+ gcc_assert (offset2 == NULL_TREE
+ || poly_int_tree_p (offset2));
tree d1 = OMP_CLAUSE_DECL (*sc);
tree d2 = OMP_CLAUSE_DECL (c);
while (TREE_CODE (d1) == ARRAY_REF)
--
2.17.1
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 658 bytes --]
^ permalink raw reply [flat|nested] 11+ messages in thread
end of thread, other threads:[~2019-11-11 9:04 UTC | newest]
Thread overview: 11+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2015-07-31 16:28 [gomp4.1] Start of structure element mapping support Jakub Jelinek
2015-08-28 18:23 ` [gomp4.1] WIP: Structure " Jakub Jelinek
2015-08-31 15:08 ` [gomp4.1] " Jakub Jelinek
2015-09-02 11:21 ` Ilya Verbin
2015-09-02 15:59 ` [gomp4.1] Depend clause support for offloading Jakub Jelinek
2015-09-03 14:19 ` Jakub Jelinek
2015-09-03 17:41 ` Jakub Jelinek
2015-09-04 10:38 ` Jakub Jelinek
2019-10-16 13:35 ` [gomp4.1] Start of structure element mapping support Thomas Schwinge
2019-10-16 17:46 ` Jakub Jelinek
2019-11-11 9:13 ` Thomas Schwinge
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).