Hi!
As the testcase shows, we weren't handling properly the case where the
decl after which .field appears in map clauses is reference to struct/class.
Fixed thusly, bootstrapped/regtested on x86_64-linux and i686-linux, and
tested with x86_64-intelmicemul-linux offloading on x86_64-linux, committed
to trunk.
2016-06-16 Jakub Jelinek <ja...@redhat.com>
* gimplify.c (gimplify_scan_omp_clauses): Handle COMPONENT_REFs
with base of reference to struct.
* parser.c (cp_parser_omp_var_list_no_open): Call
convert_from_reference before cp_parser_postfix_dot_deref_expression.
* semantics.c (finish_omp_clauses): Don't ICE when
processing_template_decl when checking for bitfields and unions.
Look through REFERENCE_REF_P as base of COMPONENT_REF.
* testsuite/libgomp.c++/target-20.C: New test.
--- gcc/gimplify.c.jj 2016-06-16 15:33:37.137940504 +0200
+++ gcc/gimplify.c 2016-06-16 16:50:21.489329556 +0200
@@ -6983,6 +6983,11 @@ gimplify_scan_omp_clauses (tree *list_p,
{
while (TREE_CODE (decl) == COMPONENT_REF)
decl = TREE_OPERAND (decl, 0);
+ if (TREE_CODE (decl) == INDIRECT_REF
+ && DECL_P (TREE_OPERAND (decl, 0))
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+ == REFERENCE_TYPE))
+ decl = TREE_OPERAND (decl, 0);
}
if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
== GS_ERROR)
@@ -6998,9 +7003,11 @@ gimplify_scan_omp_clauses (tree *list_p,
break;
}
- if (TYPE_SIZE_UNIT (TREE_TYPE (decl)) == NULL
- || (TREE_CODE (TYPE_SIZE_UNIT (TREE_TYPE (decl)))
- != INTEGER_CST))
+ tree stype = TREE_TYPE (decl);
+ if (TREE_CODE (stype) == REFERENCE_TYPE)
+ stype = TREE_TYPE (stype);
+ if (TYPE_SIZE_UNIT (stype) == NULL
+ || TREE_CODE (TYPE_SIZE_UNIT (stype)) != INTEGER_CST)
{
error_at (OMP_CLAUSE_LOCATION (c),
"mapping field %qE of variable length "
@@ -7040,6 +7047,14 @@ gimplify_scan_omp_clauses (tree *list_p,
base = get_inner_reference (base, &bitsize, &bitpos, &offset,
&mode, &unsignedp, &reversep,
&volatilep, false);
+ tree orig_base = base;
+ if ((TREE_CODE (base) == INDIRECT_REF
+ || (TREE_CODE (base) == MEM_REF
+ && integer_zerop (TREE_OPERAND (base, 1))))
+ && DECL_P (TREE_OPERAND (base, 0))
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0)))
+ == REFERENCE_TYPE))
+ base = TREE_OPERAND (base, 0);
gcc_assert (base == decl
&& (offset == NULL_TREE
|| TREE_CODE (offset) == INTEGER_CST));
@@ -7053,7 +7068,10 @@ gimplify_scan_omp_clauses (tree *list_p,
tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (l, GOMP_MAP_STRUCT);
- OMP_CLAUSE_DECL (l) = decl;
+ if (orig_base != base)
+ OMP_CLAUSE_DECL (l) = unshare_expr (orig_base);
+ else
+ OMP_CLAUSE_DECL (l) = decl;
OMP_CLAUSE_SIZE (l) = size_int (1);
if (struct_map_to_clause == NULL)
struct_map_to_clause = new hash_map<tree, tree>;
@@ -7095,6 +7113,18 @@ gimplify_scan_omp_clauses (tree *list_p,
*list_p = l;
list_p = &OMP_CLAUSE_CHAIN (l);
}
+ if (orig_base != base && code == OMP_TARGET)
+ {
+ tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ enum gomp_map_kind mkind
+ = GOMP_MAP_FIRSTPRIVATE_REFERENCE;
+ OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
+ OMP_CLAUSE_DECL (c2) = decl;
+ OMP_CLAUSE_SIZE (c2) = size_zero_node;
+ OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (l);
+ OMP_CLAUSE_CHAIN (l) = c2;
+ }
flags = GOVD_MAP | GOVD_EXPLICIT;
if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr)
flags |= GOVD_SEEN;
@@ -7113,8 +7143,12 @@ 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))
+ sc = &OMP_CLAUSE_CHAIN (*osc);
+ if (*sc != c
+ && (OMP_CLAUSE_MAP_KIND (*sc)
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+ sc = &OMP_CLAUSE_CHAIN (*sc);
+ for (; *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc))
if (ptr && sc == prev_list_p)
break;
else if (TREE_CODE (OMP_CLAUSE_DECL (*sc))
@@ -7150,6 +7184,15 @@ gimplify_scan_omp_clauses (tree *list_p,
&mode, &unsignedp,
&reversep, &volatilep,
false);
+ if ((TREE_CODE (base) == INDIRECT_REF
+ || (TREE_CODE (base) == MEM_REF
+ && integer_zerop (TREE_OPERAND (base,
+ 1))))
+ && DECL_P (TREE_OPERAND (base, 0))
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (base,
+ 0)))
+ == REFERENCE_TYPE))
+ base = TREE_OPERAND (base, 0);
if (base != decl)
break;
if (scp)
--- gcc/cp/parser.c.jj 2016-06-16 15:33:36.351950469 +0200
+++ gcc/cp/parser.c 2016-06-16 16:50:21.486329596 +0200
@@ -30001,6 +30001,7 @@ cp_parser_omp_var_list_no_open (cp_parse
= cp_lexer_peek_token (parser->lexer)->location;
cp_id_kind idk = CP_ID_KIND_NONE;
cp_lexer_consume_token (parser->lexer);
+ decl = convert_from_reference (decl);
decl
= cp_parser_postfix_dot_deref_expression (parser, CPP_DOT,
decl, false,
--- gcc/cp/semantics.c.jj 2016-06-16 15:33:36.300951116 +0200
+++ gcc/cp/semantics.c 2016-06-16 17:29:53.458543027 +0200
@@ -6664,7 +6664,8 @@ finish_omp_clauses (tree clauses, enum c
{
if (type_dependent_expression_p (t))
break;
- if (DECL_BIT_FIELD (TREE_OPERAND (t, 1)))
+ if (TREE_CODE (TREE_OPERAND (t, 1)) == FIELD_DECL
+ && DECL_BIT_FIELD (TREE_OPERAND (t, 1)))
{
error_at (OMP_CLAUSE_LOCATION (c),
"bit-field %qE in %qs clause",
@@ -6680,8 +6681,9 @@ finish_omp_clauses (tree clauses, enum c
}
while (TREE_CODE (t) == COMPONENT_REF)
{
- if (TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0)))
- == UNION_TYPE)
+ if (TREE_TYPE (TREE_OPERAND (t, 0))
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0)))
+ == UNION_TYPE))
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qE is a member of a union", t);
@@ -6692,6 +6694,8 @@ finish_omp_clauses (tree clauses, enum c
}
if (remove)
break;
+ if (REFERENCE_REF_P (t))
+ t = TREE_OPERAND (t, 0);
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
{
if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
--- libgomp/testsuite/libgomp.c++/target-20.C.jj 2016-06-16
16:48:56.172447009 +0200
+++ libgomp/testsuite/libgomp.c++/target-20.C 2016-06-16 18:27:03.783285764
+0200
@@ -0,0 +1,80 @@
+extern "C" void abort ();
+struct S { int a, b, c, d; };
+
+void
+foo (S &s)
+{
+ int err;
+ #pragma omp target map (s.b, s.d) map (from: err)
+ {
+ err = s.b != 21 || s.d != 24;
+ s.b++; s.d++;
+ }
+ if (err || s.b != 22 || s.d != 25)
+ abort ();
+ #pragma omp target data map (s.b, s.d)
+ {
+ #pragma omp target map (alloc: s.b, s.d) map (from: err)
+ {
+ err = s.b != 22 || s.d != 25;
+ s.b++; s.d++;
+ }
+ }
+ if (err || s.b != 23 || s.d != 26)
+ abort ();
+ #pragma omp target data map (s)
+ {
+ #pragma omp target map (alloc: s.b, s.d) map (from: err)
+ {
+ err = s.b != 23 || s.d != 26;
+ s.b++; s.d++;
+ }
+ }
+ if (err || s.b != 24 || s.d != 27)
+ abort ();
+}
+
+template <typename T, typename U>
+void
+bar (S &s, T &t, U u)
+{
+ int err;
+ #pragma omp target map (s.b, s.d, t.b, t.d, u.b, u.d) map (from: err)
+ {
+ err = s.b != 21 || s.d != 24 || t.b != 73 || t.d != 82 || u.b != 31 || u.d
!= 37;
+ s.b++; s.d++; t.b++; t.d++; u.b++; u.d++;
+ }
+ if (err || s.b != 22 || s.d != 25 || t.b != 74 || t.d != 83 || u.b != 32 ||
u.d != 38)
+ abort ();
+ #pragma omp target data map (s.b, s.d, t.b, t.d, u.b, u.d)
+ {
+ #pragma omp target map (alloc: s.b, s.d, t.b, t.d, u.b, u.d) map (from:
err)
+ {
+ err = s.b != 22 || s.d != 25 || t.b != 74 || t.d != 83 || u.b != 32 ||
u.d != 38;
+ s.b++; s.d++; t.b++; t.d++; u.b++; u.d++;
+ }
+ }
+ if (err || s.b != 23 || s.d != 26 || t.b != 75 || t.d != 84 || u.b != 33 ||
u.d != 39)
+ abort ();
+ #pragma omp target data map (s, t, u)
+ {
+ #pragma omp target map (alloc: s.b, s.d, t.b, t.d, u.b, u.d) map (from:
err)
+ {
+ err = s.b != 23 || s.d != 26 || t.b != 75 || t.d != 84 || u.b != 33 ||
u.d != 39;
+ s.b++; s.d++; t.b++; t.d++; u.b++; u.d++;
+ }
+ }
+ if (err || s.b != 24 || s.d != 27 || t.b != 76 || t.d != 85 || u.b != 34 ||
u.d != 40)
+ abort ();
+}
+
+int
+main ()
+{
+ S s = { 1, 21, 2, 24 };
+ foo (s);
+ S s2 = { 3, 21, 4, 24 };
+ S t = { 5, 73, 6, 82 };
+ S u = { 7, 31, 8, 37 };
+ bar <S, S &> (s2, t, u);
+}
Jakub
