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

From-SVN: r272444

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index cbf6915c828660f28e3ec8c0e9425cc013674c62..43a0a232dc2188682dbc634c98b9da322993a1df 100644 (file)
--- 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 9df21a4d0466f27771f8222293ff5f665cd88312..b0f1d94abf73f85bf157bb88e4f764720cb82d97 100644 (file)
--- 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 2848d2ceecab8326c82ad9f76a2d2cd670f745d1..473fd66d39fd5a9d115ffe68c5b8878850840143 100644 (file)
--- 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 35b1ccd367bdd516bd27f1ebba8e38914dc92d60..7c4380f4f04191cad106dab930583c6d49cc3eb1 100644 (file)
--- 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 33b82459bfc5029ba91546c9b210ed004d974a67..af43b6bc816261d49d6944b02ed139b324d2fa08 100644 (file)
--- 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 827bab2d89618ffc2c6825060876b4920377dd08..06004aafde989d6ed4e52515421bb4328c1220bd 100644 (file)
--- 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 bc726174252d4c0cedd0ecfbdbd759ab52d2d811..087b95456926c1928f9560f8df40f912182b5194 100644 (file)
--- 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;
 }