diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 9c72ecbaf5a7b10efb9a9f09f5227a15a2981c46..67cee6b5a0c48dd0d86f2e3d9666717beb09d217 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,12 @@
+2016-03-02 Jakub Jelinek <jakub@redhat.com>
+
+ PR libgomp/69555
+ * gimplify.c (gimplify_decl_expr): For decls with REFERENCE_TYPE, also
+ gimplify_type_sizes the type they refer to.
+ (omp_notice_variable): Handle reference vars to VLAs.
+ * omp-low.c (lower_omp_target): Emit setup of OMP_CLAUSE_PRIVATE reference
+ to VLA decls in the second pass instead of first pass.
+
2016-03-02 Tom de Vries <tom@codesourcery.com>
PR tree-optimization/68659
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index e7ea974496dd4db4994c1acee630de7ffb4fe6e3..b331e410a3c152a699a4a61d19feb6f97ba925c1 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -1436,14 +1436,22 @@ gimplify_decl_expr (tree *stmt_p, gimple_seq *seq_p)
if ((TREE_CODE (decl) == TYPE_DECL
|| TREE_CODE (decl) == VAR_DECL)
&& !TYPE_SIZES_GIMPLIFIED (TREE_TYPE (decl)))
- gimplify_type_sizes (TREE_TYPE (decl), seq_p);
+ {
+ gimplify_type_sizes (TREE_TYPE (decl), seq_p);
+ if (TREE_CODE (TREE_TYPE (decl)) == REFERENCE_TYPE)
+ gimplify_type_sizes (TREE_TYPE (TREE_TYPE (decl)), seq_p);
+ }
/* ??? DECL_ORIGINAL_TYPE is streamed for LTO so it needs to be gimplified
in case its size expressions contain problematic nodes like CALL_EXPR. */
if (TREE_CODE (decl) == TYPE_DECL
&& DECL_ORIGINAL_TYPE (decl)
&& !TYPE_SIZES_GIMPLIFIED (DECL_ORIGINAL_TYPE (decl)))
- gimplify_type_sizes (DECL_ORIGINAL_TYPE (decl), seq_p);
+ {
+ gimplify_type_sizes (DECL_ORIGINAL_TYPE (decl), seq_p);
+ if (TREE_CODE (DECL_ORIGINAL_TYPE (decl)) == REFERENCE_TYPE)
+ gimplify_type_sizes (TREE_TYPE (DECL_ORIGINAL_TYPE (decl)), seq_p);
+ }
if (TREE_CODE (decl) == VAR_DECL && !DECL_EXTERNAL (decl))
{
@@ -6264,16 +6272,30 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
if ((n->value & (GOVD_SEEN | GOVD_LOCAL)) == 0
&& (flags & (GOVD_SEEN | GOVD_LOCAL)) == GOVD_SEEN
- && DECL_SIZE (decl)
- && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
- {
- splay_tree_node n2;
- tree t = DECL_VALUE_EXPR (decl);
- gcc_assert (TREE_CODE (t) == INDIRECT_REF);
- t = TREE_OPERAND (t, 0);
- gcc_assert (DECL_P (t));
- n2 = splay_tree_lookup (ctx->variables, (splay_tree_key) t);
- n2->value |= GOVD_SEEN;
+ && DECL_SIZE (decl))
+ {
+ if (TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
+ {
+ splay_tree_node n2;
+ tree t = DECL_VALUE_EXPR (decl);
+ gcc_assert (TREE_CODE (t) == INDIRECT_REF);
+ t = TREE_OPERAND (t, 0);
+ gcc_assert (DECL_P (t));
+ n2 = splay_tree_lookup (ctx->variables, (splay_tree_key) t);
+ n2->value |= GOVD_SEEN;
+ }
+ else if (lang_hooks.decls.omp_privatize_by_reference (decl)
+ && TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (decl)))
+ && (TREE_CODE (TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (decl))))
+ != INTEGER_CST))
+ {
+ splay_tree_node n2;
+ tree t = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (decl)));
+ gcc_assert (DECL_P (t));
+ n2 = splay_tree_lookup (ctx->variables, (splay_tree_key) t);
+ if (n2)
+ n2->value |= GOVD_SEEN;
+ }
}
shared = ((flags | n->value) & GOVD_SHARED) != 0;
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 989d03ef7740a9c108fd4740e47ebb5eef1f4b0b..ecbf74ad71bf7201c0e4a5fe4aabea8d1cf895d1 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -16472,13 +16472,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
x = build_fold_addr_expr_loc (clause_loc, x);
}
else
- {
- tree atmp
- = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
- tree rtype = TREE_TYPE (TREE_TYPE (new_var));
- tree al = size_int (TYPE_ALIGN (rtype));
- x = build_call_expr_loc (clause_loc, atmp, 2, x, al);
- }
+ break;
x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
@@ -16545,7 +16539,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
}
/* Handle GOMP_MAP_FIRSTPRIVATE_{POINTER,REFERENCE} in second pass,
so that firstprivate vars holding OMP_CLAUSE_SIZE if needed
- are already handled. */
+ are already handled. Similarly OMP_CLAUSE_PRIVATE for VLAs
+ or references to VLAs. */
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c))
{
@@ -16687,6 +16682,27 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gimple_seq_add_stmt (&new_body,
gimple_build_assign (new_pvar, x));
}
+ else if (is_reference (var) && !is_gimple_omp_oacc (ctx->stmt))
+ {
+ location_t clause_loc = OMP_CLAUSE_LOCATION (c);
+ tree new_var = lookup_decl (var, ctx);
+ tree x = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (new_var)));
+ if (TREE_CONSTANT (x))
+ break;
+ else
+ {
+ tree atmp
+ = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN);
+ tree rtype = TREE_TYPE (TREE_TYPE (new_var));
+ tree al = size_int (TYPE_ALIGN (rtype));
+ x = build_call_expr_loc (clause_loc, atmp, 2, x, al);
+ }
+
+ x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
+ gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_assign (new_var, x));
+ }
break;
}
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index e96f1495fc255044b19c70e65b2952fcef6250b8..90f245dd3190031f07e9efd89742964bf0aff299 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,9 @@
+2016-03-02 Jakub Jelinek <jakub@redhat.com>
+
+ PR libgomp/69555
+ * testsuite/libgomp.c++/pr69555-1.C: New test.
+ * testsuite/libgomp.c++/pr69555-2.C: New test.
+
2016-02-26 Keith McDaniel <k.allen.mcdaniel@gmail.com>
Martin Jambor <mjambor@suse.cz>
diff --git a/libgomp/testsuite/libgomp.c++/pr69555-1.C b/libgomp/testsuite/libgomp.c++/pr69555-1.C
new file mode 100644
index 0000000000000000000000000000000000000000..c6885ff73213de77a3b2ca9e53c3e5ae8237e2d6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr69555-1.C
@@ -0,0 +1,114 @@
+// PR libgomp/69555
+// { dg-do run }
+
+#include <omp.h>
+
+__attribute__((noinline, noclone)) void
+f1 (int y)
+{
+ int a[y - 2];
+ int (&c)[y - 2] = a;
+ c[0] = 111;
+ int e = 0;
+
+ #pragma omp parallel private (c) num_threads (4) reduction (+:e)
+ {
+ int v = omp_get_thread_num ();
+ for (int i = 0; i < y - 2; i++)
+ c[i] = i + v;
+ #pragma omp barrier
+ for (int i = 0; i < y - 2; i++)
+ if (c[i] != i + v)
+ e++;
+ }
+ if (c[0] != 111 || e)
+ __builtin_abort ();
+}
+
+__attribute__((noinline, noclone)) void
+f2 (int y)
+{
+ int a[y - 2];
+ int (&c)[y - 2] = a;
+ c[0] = 111;
+
+ #pragma omp task private (c)
+ {
+ int v = omp_get_thread_num ();
+ for (int i = 0; i < y - 2; i++)
+ c[i] = i + v;
+ asm volatile ("" : : "r" (&c[0]) : "memory");
+ for (int i = 0; i < y - 2; i++)
+ if (c[i] != i + v)
+ __builtin_abort ();
+ }
+ if (c[0] != 111)
+ __builtin_abort ();
+}
+
+__attribute__((noinline, noclone)) void
+f3 (int y)
+{
+ int a[y - 2];
+ int (&c)[y - 2] = a;
+ for (int i = 0; i < y - 2; i++)
+ c[i] = i + 4;
+
+ #pragma omp parallel firstprivate (c) num_threads (4)
+ {
+ int v = omp_get_thread_num ();
+ for (int i = 0; i < y - 2; i++)
+ {
+ if (c[i] != i + 4)
+ __builtin_abort ();
+ c[i] = i + v;
+ }
+ #pragma omp barrier
+ for (int i = 0; i < y - 2; i++)
+ if (c[i] != i + v)
+ __builtin_abort ();
+ }
+ for (int i = 0; i < y - 2; i++)
+ if (c[i] != i + 4)
+ __builtin_abort ();
+}
+
+__attribute__((noinline, noclone)) void
+f4 (int y)
+{
+ int a[y - 2];
+ int (&c)[y - 2] = a;
+ for (int i = 0; i < y - 2; i++)
+ c[i] = i + 4;
+
+ #pragma omp task firstprivate (c)
+ {
+ int v = omp_get_thread_num ();
+ for (int i = 0; i < y - 2; i++)
+ {
+ if (c[i] != i + 4)
+ __builtin_abort ();
+ c[i] = i + v;
+ }
+ asm volatile ("" : : "r" (&c[0]) : "memory");
+ for (int i = 0; i < y - 2; i++)
+ if (c[i] != i + v)
+ __builtin_abort ();
+ }
+ for (int i = 0; i < y - 2; i++)
+ if (c[i] != i + 4)
+ __builtin_abort ();
+}
+
+int
+main ()
+{
+ f1 (6);
+ f3 (6);
+ #pragma omp parallel num_threads (4)
+ {
+ f2 (6);
+ f4 (6);
+ }
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/pr69555-2.C b/libgomp/testsuite/libgomp.c++/pr69555-2.C
new file mode 100644
index 0000000000000000000000000000000000000000..78776ea7c9e1fed804183f9e4217678ff447386d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr69555-2.C
@@ -0,0 +1,58 @@
+// PR libgomp/69555
+// { dg-do run }
+
+__attribute__((noinline, noclone)) void
+f1 (int y)
+{
+ int a[y - 2];
+ int (&c)[y - 2] = a;
+ for (int i = 0; i < y - 2; i++)
+ c[i] = i + 4;
+
+ #pragma omp target firstprivate (c)
+ {
+ for (int i = 0; i < y - 2; i++)
+ {
+ if (c[i] != i + 4)
+ __builtin_abort ();
+ c[i] = i + 9;
+ }
+ asm volatile ("" : : "r" (&c[0]) : "memory");
+ for (int i = 0; i < y - 2; i++)
+ if (c[i] != i + 9)
+ __builtin_abort ();
+ }
+ for (int i = 0; i < y - 2; i++)
+ if (c[i] != i + 4)
+ __builtin_abort ();
+}
+
+__attribute__((noinline, noclone)) void
+f2 (int y)
+{
+ int a[y - 2];
+ int (&c)[y - 2] = a;
+ for (int i = 0; i < y - 2; i++)
+ c[i] = i + 4;
+
+ #pragma omp target private (c)
+ {
+ for (int i = 0; i < y - 2; i++)
+ c[i] = i + 9;
+ asm volatile ("" : : "r" (&c[0]) : "memory");
+ for (int i = 0; i < y - 2; i++)
+ if (c[i] != i + 9)
+ __builtin_abort ();
+ }
+ for (int i = 0; i < y - 2; i++)
+ if (c[i] != i + 4)
+ __builtin_abort ();
+}
+
+int
+main ()
+{
+ f1 (6);
+ f2 (6);
+ return 0;
+}