From acb4157074770f715968c3a9c1e6929f98fcddc8 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Tue, 18 Jun 2019 22:13:54 +0000
Subject: [PATCH] [PR90862] OpenACC 'declare' ICE when nested inside another
 construct

	gcc/
	PR middle-end/90862
	* omp-low.c (check_omp_nesting_restrictions): Handle
	GF_OMP_TARGET_KIND_OACC_DECLARE.
	gcc/testsuite/
	PR middle-end/90862
	* c-c++-common/goacc/declare-1.c: Update.
	* c-c++-common/goacc/declare-2.c: Likewise.
	libgomp/
	PR middle-end/90862
	* testsuite/libgomp.oacc-c-c++-common/declare-1.c: Update.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@272444 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/ChangeLog                                 |  6 ++
 gcc/omp-low.c                                 |  1 +
 gcc/testsuite/ChangeLog                       |  6 ++
 gcc/testsuite/c-c++-common/goacc/declare-1.c  | 82 +++++++++++++++-
 gcc/testsuite/c-c++-common/goacc/declare-2.c  | 35 ++++++-
 libgomp/ChangeLog                             |  5 +
 .../libgomp.oacc-c-c++-common/declare-1.c     | 98 +++++++++++++++++--
 7 files changed, 223 insertions(+), 10 deletions(-)

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index cbf6915c8286..43a0a232dc21 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,9 @@
+2019-06-18  Thomas Schwinge  <thomas@codesourcery.com>
+
+	PR middle-end/90862
+	* omp-low.c (check_omp_nesting_restrictions): Handle
+	GF_OMP_TARGET_KIND_OACC_DECLARE.
+
 2019-06-18  Uroš Bizjak  <ubizjak@gmail.com>
 
 	* config/i386/i386.md (@cmp<mode>_1): Rename from cmp<mode>_1.
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 9df21a4d0466..b0f1d94abf73 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -3119,6 +3119,7 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
 	    case GF_OMP_TARGET_KIND_OACC_UPDATE: stmt_name = "update"; break;
 	    case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
 	      stmt_name = "enter/exit data"; break;
+	    case GF_OMP_TARGET_KIND_OACC_DECLARE: stmt_name = "declare"; break;
 	    case GF_OMP_TARGET_KIND_OACC_HOST_DATA: stmt_name = "host_data";
 	      break;
 	    default: gcc_unreachable ();
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 2848d2ceecab..473fd66d39fd 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,9 @@
+2019-06-18  Thomas Schwinge  <thomas@codesourcery.com>
+
+	PR middle-end/90862
+	* c-c++-common/goacc/declare-1.c: Update.
+	* c-c++-common/goacc/declare-2.c: Likewise.
+
 2019-06-18  Marek Polacek  <polacek@redhat.com>
 
 	PR c++/84698
diff --git a/gcc/testsuite/c-c++-common/goacc/declare-1.c b/gcc/testsuite/c-c++-common/goacc/declare-1.c
index 35b1ccd367bd..7c4380f4f041 100644
--- a/gcc/testsuite/c-c++-common/goacc/declare-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/declare-1.c
@@ -1,5 +1,5 @@
-/* Test valid uses of declare directive.  */
-/* { dg-do compile } */
+/* Test valid use of the OpenACC 'declare' directive.  */
+
 
 int v0;
 #pragma acc declare create(v0)
@@ -25,6 +25,7 @@ int v9;
 int v10;
 #pragma acc declare present_or_create(v10)
 
+
 void
 f (void)
 {
@@ -93,3 +94,80 @@ f (void)
   }
  b:;
 }
+
+
+/* The same as 'f' but everything contained in an OpenACC 'data' construct.  */
+
+void
+f_data (void)
+{
+#pragma acc data
+  {
+    int va0;
+# pragma acc declare create(va0)
+
+    int va1;
+# pragma acc declare copyin(va1)
+
+    int *va2;
+# pragma acc declare deviceptr(va2)
+
+    int va3;
+# pragma acc declare device_resident(va3)
+
+#if 0 /* TODO */
+    extern int ve0;
+# pragma acc declare create(ve0)
+
+    extern int ve1;
+# pragma acc declare copyin(ve1)
+
+    extern int *ve2;
+# pragma acc declare deviceptr(ve2)
+
+    extern int ve3;
+# pragma acc declare device_resident(ve3)
+
+    extern int ve4;
+# pragma acc declare link(ve4)
+
+    extern int ve5;
+# pragma acc declare present_or_copyin(ve5)
+ 
+    extern int ve6;
+# pragma acc declare present_or_create(ve6)
+#endif
+
+    int va5;
+# pragma acc declare copy(va5)
+
+    int va6;
+# pragma acc declare copyout(va6)
+
+    int va7;
+# pragma acc declare present(va7)
+
+    int va8;
+# pragma acc declare present_or_copy(va8)
+
+    int va9;
+# pragma acc declare present_or_copyin(va9)
+
+    int va10;
+# pragma acc declare present_or_copyout(va10)
+
+    int va11;
+# pragma acc declare present_or_create(va11)
+
+  a:
+    {
+      int va0;
+# pragma acc declare create(va0)
+      if (v1)
+	goto a;
+      else
+	goto b;
+    }
+  b:;
+  }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/declare-2.c b/gcc/testsuite/c-c++-common/goacc/declare-2.c
index 33b82459bfc5..af43b6bc8162 100644
--- a/gcc/testsuite/c-c++-common/goacc/declare-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/declare-2.c
@@ -1,5 +1,5 @@
-/* Test invalid uses of declare directive.  */
-/* { dg-do compile } */
+/* Test invalid use of the OpenACC 'declare' directive.  */
+
 
 #pragma acc declare /* { dg-error "no valid clauses" } */
 
@@ -42,6 +42,7 @@ int va11;
 int va12;
 #pragma acc declare create (va12) link (va12) /* { dg-error "more than once" } */
 
+
 void
 f (void)
 {
@@ -65,3 +66,33 @@ f (void)
 
 #pragma acc declare present (v2) /* { dg-error "invalid use of" } */
 }
+
+
+/* The same as 'f' but everything contained in an OpenACC 'data' construct.  */
+
+void
+f_data (void)
+{
+#pragma acc data
+  {
+    int va0;
+# pragma acc declare link(va0) /* { dg-error "global variable" } */
+
+    extern int ve0;
+# pragma acc declare copy(ve0) /* { dg-error "invalid use of" } */
+
+    extern int ve1;
+# pragma acc declare copyout(ve1) /* { dg-error "invalid use of" } */
+
+    extern int ve2;
+# pragma acc declare present(ve2) /* { dg-error "invalid use of" } */
+
+    extern int ve3;
+# pragma acc declare present_or_copy(ve3) /* { dg-error "invalid use of" } */
+
+    extern int ve4;
+# pragma acc declare present_or_copyout(ve4) /* { dg-error "invalid use of" } */
+
+# pragma acc declare present (v2) /* { dg-error "invalid use of" } */
+  }
+}
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 827bab2d8961..06004aafde98 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,8 @@
+2019-06-18  Thomas Schwinge  <thomas@codesourcery.com>
+
+	PR middle-end/90862
+	* testsuite/libgomp.oacc-c-c++-common/declare-1.c: Update.
+
 2019-06-16  Tom de Vries  <tdevries@suse.de>
 
 	PR tree-optimization/89376
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c
index bc726174252d..087b95456926 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c
@@ -1,6 +1,5 @@
 #include <openacc.h>
 #include <stdlib.h>
-#include <stdio.h>
 
 #define N 8
 
@@ -39,14 +38,14 @@ subr1 (int *a)
   }
 }
 
-int b[8];
+int b[N];
 #pragma acc declare create (b)
 
-int d[8] = { 1, 2, 3, 4, 5, 6, 7, 8 };
+int d[N] = { 1, 2, 3, 4, 5, 6, 7, 8 };
 #pragma acc declare copyin (d)
 
-int
-main (int argc, char **argv)
+static void
+f (void)
 {
   int a[N];
   int e[N];
@@ -110,11 +109,98 @@ main (int argc, char **argv)
 
   subr2 (&a[0]);
 
-  for (i = 0; i < 1; i++)
+  for (i = 0; i < N; i++)
     {
       if (a[i] != 1234 * 6)
 	abort ();
     }
+}
+
+
+/* The same as 'f' but everything contained in an OpenACC 'data' construct.  */
+
+static void
+f_data (void)
+{
+#pragma acc data
+  {
+    int a[N];
+    int e[N];
+# pragma acc declare create (e)
+    int i;
+
+    for (i = 0; i < N; i++)
+      a[i] = i + 1;
+
+    if (!acc_is_present (&b, sizeof (b)))
+      abort ();
+
+    if (!acc_is_present (&d, sizeof (d)))
+      abort ();
+
+    if (!acc_is_present (&e, sizeof (e)))
+      abort ();
+
+# pragma acc parallel copyin (a[0:N])
+    {
+      for (i = 0; i < N; i++)
+	{
+	  b[i] = a[i];
+	  a[i] = b[i];
+	}
+    }
+
+    for (i = 0; i < N; i++)
+      {
+	if (a[i] != i + 1)
+	  abort ();
+      }
+
+# pragma acc parallel copy (a[0:N])
+    {
+      for (i = 0; i < N; i++)
+	{
+	  e[i] = a[i] + d[i];
+	  a[i] = e[i];
+	}
+    }
+
+    for (i = 0; i < N; i++)
+      {
+	if (a[i] != (i + 1) * 2)
+	  abort ();
+      }
+
+    for (i = 0; i < N; i++)
+      {
+	a[i] = 1234;
+      }
+
+    subr1 (&a[0]);
+
+    for (i = 0; i < N; i++)
+      {
+	if (a[i] != 1234 * 2)
+	  abort ();
+      }
+
+    subr2 (&a[0]);
+
+    for (i = 0; i < N; i++)
+      {
+	if (a[i] != 1234 * 6)
+	  abort ();
+      }
+  }
+}
+
+
+int
+main (int argc, char **argv)
+{
+  f ();
+
+  f_data ();
 
   return 0;
 }
-- 
2.20.1

