diff options
| -rw-r--r-- | gcc/ChangeLog | 6 | ||||
| -rw-r--r-- | gcc/omp-low.c | 1 | ||||
| -rw-r--r-- | gcc/testsuite/ChangeLog | 6 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/declare-1.c | 82 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/declare-2.c | 35 | ||||
| -rw-r--r-- | libgomp/ChangeLog | 5 | ||||
| -rw-r--r-- | libgomp/testsuite/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 cbf6915..43a0a23 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 9df21a4..b0f1d94 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 2848d2c..473fd66 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 35b1ccd..7c4380f 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 33b8245..af43b6b 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 827bab2..06004aa 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 bc7261742..087b954 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; } |