In addition to resolving the memory leak involving OpenACC delare
clauses, this patch also corrects an ICE involving VLA variables as data
clause arguments used in acc declare constructs. More details can be
found here <https://gcc.gnu.org/ml/gcc-patches/2016-09/msg01630.html>.
Is this OK for trunk?
Cesar
2017-03-22 Cesar Philippidis <ce...@codesourcery.com>
PR c++/80029
gcc/
* gimplify.c (is_oacc_declared): New function.
(oacc_default_clause): Use it to set default flags for acc declared
variables inside parallel regions.
(gimplify_scan_omp_clauses): Strip firstprivate pointers for acc
declared variables.
(gimplify_oacc_declare): Gimplify the declare clauses. Add the
declare attribute to any decl as necessary.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/declare-vla.c: New test.
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index fbf136f..26d35d1 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -6787,6 +6787,16 @@ device_resident_p (tree decl)
return false;
}
+/* Return true if DECL has an ACC DECLARE attribute. */
+
+static bool
+is_oacc_declared (tree decl)
+{
+ tree t = TREE_CODE (decl) == MEM_REF ? TREE_OPERAND (decl, 0) : decl;
+ tree declared = lookup_attribute ("oacc declare target", DECL_ATTRIBUTES (t));
+ return declared != NULL_TREE;
+}
+
/* Determine outer default flags for DECL mentioned in an OMP region
but not declared in an enclosing clause.
@@ -6887,6 +6897,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
{
const char *rkind;
bool on_device = false;
+ bool declared = is_oacc_declared (decl);
tree type = TREE_TYPE (decl);
if (lang_hooks.decls.omp_privatize_by_reference (decl))
@@ -6917,7 +6928,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
case ORT_ACC_PARALLEL:
{
- if (on_device || AGGREGATE_TYPE_P (type))
+ if (on_device || AGGREGATE_TYPE_P (type) || declared)
/* Aggregates default to 'present_or_copy'. */
flags |= GOVD_MAP;
else
@@ -7346,6 +7357,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_TARGET_DATA:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
+ case OACC_DECLARE:
case OACC_HOST_DATA:
ctx->target_firstprivatize_array_bases = true;
default:
@@ -9231,18 +9243,26 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p)
{
tree expr = *expr_p;
gomp_target *stmt;
- tree clauses, t;
+ tree clauses, t, decl;
clauses = OACC_DECLARE_CLAUSES (expr);
gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA, OACC_DECLARE);
+ gimplify_adjust_omp_clauses (pre_p, NULL, &clauses, OACC_DECLARE);
for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
{
- tree decl = OMP_CLAUSE_DECL (t);
+ decl = OMP_CLAUSE_DECL (t);
if (TREE_CODE (decl) == MEM_REF)
- continue;
+ decl = TREE_OPERAND (decl, 0);
+
+ if (VAR_P (decl) && !is_oacc_declared (decl))
+ {
+ tree attr = get_identifier ("oacc declare target");
+ DECL_ATTRIBUTES (decl) = tree_cons (attr, NULL_TREE,
+ DECL_ATTRIBUTES (decl));
+ }
if (VAR_P (decl)
&& !is_global_var (decl)
@@ -9258,7 +9278,8 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p)
}
}
- omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN);
+ if (gimplify_omp_ctxp)
+ omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN);
}
stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE,
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c
new file mode 100644
index 0000000..3ea148e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c
@@ -0,0 +1,25 @@
+/* Verify that acc declare accept VLA variables. */
+
+#include <assert.h>
+
+int
+main ()
+{
+ int N = 1000;
+ int i, A[N];
+#pragma acc declare copy(A)
+
+ for (i = 0; i < N; i++)
+ A[i] = -i;
+
+#pragma acc kernels
+ for (i = 0; i < N; i++)
+ A[i] = i;
+
+#pragma acc update host(A)
+
+ for (i = 0; i < N; i++)
+ assert (A[i] == i);
+
+ return 0;
+}
