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(-)
@@ -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.
@@ -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 ();
@@ -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
@@ -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:;
+ }
+}
@@ -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" } */
+ }
+}
@@ -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
@@ -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