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.


git-svn-id: https://gcc.gnu.org/svn/gcc/trunk@246381 138bc75d-0d04-0410-961f-82ee72b054a4

4 files changed, 68 insertions, 6 deletions

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 18d00c29fba..f7202cf0b90 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,4 +1,15 @@
-t2017-03-22  Thomas Preud'homme  <thomas.preudhomme@arm.com>
+2017-03-22  Cesar Philippidis  <cesar@codesourcery.com>
+
+	PR c++/80029
+	* 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.
+
+2017-03-22  Thomas Preud'homme  <thomas.preudhomme@arm.com>
 
 	PR target/80082
 	* config/arm/arm-isa.h (isa_bit_lpae): New feature bit.
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 06f984cec01..5658d0a2f0b 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -6786,6 +6786,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.
 
@@ -6886,6 +6896,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))
@@ -6916,7 +6927,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
@@ -7345,6 +7356,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:
@@ -9230,18 +9242,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)
@@ -9257,7 +9277,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/ChangeLog b/libgomp/ChangeLog
index f63f028d3e5..74f50e0f32f 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,8 @@
+2017-03-22  Cesar Philippidis  <cesar@codesourcery.com>
+
+	PR c++/80029
+	* testsuite/libgomp.oacc-c-c++-common/declare-vla.c: New test.
+
 2017-03-08  Jakub Jelinek  <jakub@redhat.com>
 
 	PR c/79940
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 00000000000..3ea148ed40d
--- /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;
+}