[openacc] fix an ICE with acc declared VLAs

[openacc] fix an ICE with acc declared VLAs

[Cesar Philippidis]

[-- Attachment #1: Type: text/plain, Size: 1151 bytes --]

There's a bug with ACC DECLARE clauses involving VLA variables that
causes lower_omp_target to thow an ICE on VLA decls. The problem
occurred because those clauses are never gimplified. This patch resolves
that ICE by teaching gimplify_oacc_declare how to gimplify them.

It turns out that gimplify_adjust_omp_clauses removes the
gimplify_omp_ctx created by gimplify_scan_omp_clauses. To workaround
that problem, I just taught gimplify_oacc_declare to ensure that each
gimplified decl has an "oacc declare target" attribute. This seems like
a reasonable solution because acc declare is an executable directive,
and hence really doesn't have a normal scope like ACC DATA. The declare
code does install destructors to unmap those clauses (except for those
vars which were ACC DECLARE CREATE'd -- but that's a problem for patch),
but that doesn't depend on the gimplify omp ctx. The only other major
thing that had to be updated was oacc_default_clause, because
omp_notice_variable no longer has a gimple omp ctx to inspect for
existing variables. But that's where the "oacc declare target" attribute
comes into play.

Is this patch OK for trunk?

Cesar

[-- Attachment #2: trunk-declare_copy.diff --]
[-- Type: text/x-patch, Size: 3928 bytes --]

2016-09-22  Cesar Philippidis  <cesar@codesourcery.com>

	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 e378ed0..19095ff 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -6114,6 +6114,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.
 
@@ -6214,6 +6224,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))
@@ -6244,7 +6255,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
@@ -6714,6 +6725,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:
@@ -8589,20 +8601,28 @@ 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 (TREE_CODE (decl) == VAR_DECL
+      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)
 	  && DECL_CONTEXT (decl) == current_function_decl)
 	{
@@ -8616,7 +8636,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;
+}

[gomp4] fix an ICE with acc declared VLAs

[Cesar Philippidis]

[-- Attachment #1: Type: text/plain, Size: 239 bytes --]

I've applied this patch to gomp4 which fixes a problem with VLA
variables in data clauses used in acc declare directives. More details
regarding this fix can be found here
<https://gcc.gnu.org/ml/gcc-patches/2016-09/msg01630.html>.

Cesar

[-- Attachment #2: gomp4-declare_copy.diff --]
[-- Type: text/x-patch, Size: 3960 bytes --]

2016-09-23  Cesar Philippidis  <cesar@codesourcery.com>

	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 1ecfaaa..1c04b33 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -6002,6 +6002,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.
 
@@ -6103,6 +6113,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
   const char *rkind;
   bool on_device = false;
   bool is_private = false;
+  bool declared = is_oacc_declared (decl);
   tree type = TREE_TYPE (decl);
 
   if (lang_hooks.decls.omp_privatize_by_reference (decl))
@@ -6137,7 +6148,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
 
     case ORT_ACC_PARALLEL:
       {
-	if (!is_private && (on_device || AGGREGATE_TYPE_P (type)))
+	if (!is_private && (on_device || AGGREGATE_TYPE_P (type) || declared))
 	  /* Aggregates default to 'present_or_copy'.  */
 	  flags |= GOVD_MAP;
 	else
@@ -6613,6 +6624,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
       case OMP_TARGET_ENTER_DATA:
       case OMP_TARGET_EXIT_DATA:
 	//case OACC_DATA:
+      case OACC_DECLARE:
       case OACC_HOST_DATA:
 	//case OACC_PARALLEL:
 	//case OACC_KERNELS:
@@ -8538,20 +8550,28 @@ 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 (TREE_CODE (decl) == VAR_DECL
+      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)
 	  && DECL_CONTEXT (decl) == current_function_decl)
 	{
@@ -8565,7 +8585,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;
+}

[PATCH] Fix PR80029

[Cesar Philippidis]

[-- Attachment #1: Type: text/plain, Size: 309 bytes --]

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

[-- Attachment #2: pr80029.diff --]
[-- Type: text/x-patch, Size: 3846 bytes --]

2017-03-22  Cesar Philippidis  <cesar@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;
+}

Re: [PATCH] Fix PR80029

[Jakub Jelinek]

On Wed, Mar 22, 2017 at 06:40:03AM -0700, Cesar Philippidis wrote:
> 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  <cesar@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.

Ok if testing passed.

	Jakub

Re: [PATCH] Fix PR80029

[Cesar Philippidis]

On 03/22/2017 06:42 AM, Jakub Jelinek wrote:

>> 2017-03-22  Cesar Philippidis  <cesar@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.
> 
> Ok if testing passed.

Yes it did. By the way, is there a way to automate valgrind testing, or
is that something that require manual inspection?

Cesar

Re: [PATCH] Fix PR80029

[Jakub Jelinek]

On Wed, Mar 22, 2017 at 06:48:09AM -0700, Cesar Philippidis wrote:
> On 03/22/2017 06:42 AM, Jakub Jelinek wrote:
> 
> >> 2017-03-22  Cesar Philippidis  <cesar@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.
> > 
> > Ok if testing passed.
> 
> Yes it did. By the way, is there a way to automate valgrind testing, or
> is that something that require manual inspection?

--enable-checking=valgrind if you have enough CPU time to burn (like a
weekend for a single bootstrap+regtest).

	Jakub

Re: [PATCH] Fix PR80029

[Thomas Schwinge]

Hi Cesar!

On Wed, 22 Mar 2017 06:40:03 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> 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?

(Got committed in r246381.)

> 	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.

> --- 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)
> +{

Adding here a "gcc_assert (TREE_CODE (decl) != MEM_REF);", that one never
triggers (at least given the current test coverage), so...

> +  tree t = TREE_CODE (decl) == MEM_REF ? TREE_OPERAND (decl, 0) : decl;

... it would seem that this code can be simplified -- or under which
circumstances should we get a MEM_REF here?  The call from
gimplify_oacc_declare itself does the tree operand 0 indirection, and the
call from oacc_default_clause shouldn't be a MEM_REF, I would think.

> +  tree declared = lookup_attribute ("oacc declare target", DECL_ATTRIBUTES (t));
> +  return declared != NULL_TREE;
> +}

> @@ -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

Isn't this new "declared" logic doing a thing very similar to the
existing "on_device" logic?  Should that be combined/cleaned up?

Why doesn't the same "declared"/"on_device" logic also apply to
ORT_ACC_KERNELS?


Grüße
 Thomas


> @@ -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;
> +}