This patch updates the existing OpenACC compiler tests with the new OpenACC 2.5 data clause semantics.
Is it OK for trunk? Cesar
2018-06-19 Chung-Lin Tang <clt...@codesourcery.com> Thomas Schwinge <tho...@codesourcery.com> Cesar Philippidis <ce...@codesourcery.com> gcc/testsuite/ * c-c++-common/goacc/declare-1.c: Update test case to utilize OpenACC 2.5 data clause semantics. * c-c++-common/goacc/declare-2.c: Likewise. * c-c++-common/goacc/default-4.c: Likewise. * c-c++-common/goacc/finalize-1.c: New test. * c-c++-common/goacc/kernels-alias-2.c: Update test case to utilize OpenACC 2.5 data clause semantics. * c-c++-common/goacc/kernels-alias.c: Likewise. * c-c++-common/goacc/routine-5.c: Likewise. * c-c++-common/goacc/update-if_present-1.c: New test. * c-c++-common/goacc/update-if_present-2.c: New test. * g++.dg/goacc/template.C: Update test case to utilize OpenACC 2.5 data clause semantics. * gfortran.dg/goacc/combined-directives.f90: Likewise. * gfortran.dg/goacc/data-tree.f95: Likewise. * gfortran.dg/goacc/declare-2.f95: Likewise. * gfortran.dg/goacc/default-4.f: Likewise. * gfortran.dg/goacc/enter-exit-data.f95: Likewise. * gfortran.dg/goacc/finalize-1.f: New test. * gfortran.dg/goacc/kernels-alias-2.f95: Update test case to utilize OpenACC 2.5 data clause semantics. * gfortran.dg/goacc/kernels-alias.f95: Likewise. * gfortran.dg/goacc/kernels-tree.f95: Likewise. * gfortran.dg/goacc/nested-function-1.f90: Likewise. * gfortran.dg/goacc/parallel-tree.f95: Likewise. * gfortran.dg/goacc/reduction-promotions.f90: Likewise. * gfortran.dg/goacc/update-if_present-1.f90: New test. * gfortran.dg/goacc/update-if_present-2.f90: New test. >From 87819f5846a3d4aae3983740e09a0ac4e1eb866f Mon Sep 17 00:00:00 2001 From: Cesar Philippidis <ce...@codesourcery.com> Date: Tue, 19 Jun 2018 09:28:54 -0700 Subject: [PATCH 1/7] compiler tests --- gcc/testsuite/c-c++-common/goacc/declare-1.c | 12 ++ gcc/testsuite/c-c++-common/goacc/declare-2.c | 18 +-- gcc/testsuite/c-c++-common/goacc/default-4.c | 6 +- gcc/testsuite/c-c++-common/goacc/finalize-1.c | 28 ++++ .../c-c++-common/goacc/kernels-alias-2.c | 10 +- .../c-c++-common/goacc/kernels-alias.c | 10 +- gcc/testsuite/c-c++-common/goacc/routine-5.c | 150 +++++++----------- .../c-c++-common/goacc/update-if_present-1.c | 28 ++++ .../c-c++-common/goacc/update-if_present-2.c | 42 +++++ gcc/testsuite/g++.dg/goacc/template.C | 13 +- .../gfortran.dg/goacc/combined-directives.f90 | 2 +- gcc/testsuite/gfortran.dg/goacc/data-tree.f95 | 8 +- gcc/testsuite/gfortran.dg/goacc/declare-2.f95 | 6 +- gcc/testsuite/gfortran.dg/goacc/default-4.f | 6 +- .../gfortran.dg/goacc/enter-exit-data.f95 | 3 + gcc/testsuite/gfortran.dg/goacc/finalize-1.f | 27 ++++ .../gfortran.dg/goacc/kernels-alias-2.f95 | 10 +- .../gfortran.dg/goacc/kernels-alias.f95 | 10 +- .../gfortran.dg/goacc/kernels-tree.f95 | 8 +- .../gfortran.dg/goacc/nested-function-1.f90 | 8 + .../gfortran.dg/goacc/parallel-tree.f95 | 12 +- .../goacc/reduction-promotions.f90 | 6 +- .../gfortran.dg/goacc/update-if_present-1.f90 | 27 ++++ .../gfortran.dg/goacc/update-if_present-2.f90 | 52 ++++++ 24 files changed, 345 insertions(+), 157 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/finalize-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/update-if_present-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/update-if_present-2.c create mode 100644 gcc/testsuite/gfortran.dg/goacc/finalize-1.f create mode 100644 gcc/testsuite/gfortran.dg/goacc/update-if_present-1.f90 create mode 100644 gcc/testsuite/gfortran.dg/goacc/update-if_present-2.f90 diff --git a/gcc/testsuite/c-c++-common/goacc/declare-1.c b/gcc/testsuite/c-c++-common/goacc/declare-1.c index b036c636166..35b1ccd367b 100644 --- a/gcc/testsuite/c-c++-common/goacc/declare-1.c +++ b/gcc/testsuite/c-c++-common/goacc/declare-1.c @@ -19,6 +19,12 @@ int v4; int v5, v6, v7, v8; #pragma acc declare create(v5, v6) copyin(v7, v8) +int v9; +#pragma acc declare present_or_copyin(v9) + +int v10; +#pragma acc declare present_or_create(v10) + void f (void) { @@ -49,6 +55,12 @@ f (void) 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) + int va5; #pragma acc declare copy(va5) diff --git a/gcc/testsuite/c-c++-common/goacc/declare-2.c b/gcc/testsuite/c-c++-common/goacc/declare-2.c index e41a0f59537..33b82459bfc 100644 --- a/gcc/testsuite/c-c++-common/goacc/declare-2.c +++ b/gcc/testsuite/c-c++-common/goacc/declare-2.c @@ -29,13 +29,7 @@ int v6; #pragma acc declare present_or_copy(v6) /* { dg-error "at file scope" } */ int v7; -#pragma acc declare present_or_copyin(v7) /* { dg-error "at file scope" } */ - -int v8; -#pragma acc declare present_or_copyout(v8) /* { dg-error "at file scope" } */ - -int v9; -#pragma acc declare present_or_create(v9) /* { dg-error "at file scope" } */ +#pragma acc declare present_or_copyout(v7) /* { dg-error "at file scope" } */ int va10; #pragma acc declare create (va10) @@ -67,13 +61,7 @@ f (void) #pragma acc declare present_or_copy(ve3) /* { dg-error "invalid use of" } */ extern int ve4; -#pragma acc declare present_or_copyin(ve4) /* { dg-error "invalid use of" } */ - - extern int ve5; -#pragma acc declare present_or_copyout(ve5) /* { dg-error "invalid use of" } */ - - extern int ve6; -#pragma acc declare present_or_create(ve6) /* { dg-error "invalid use of" } */ +#pragma acc declare present_or_copyout(ve4) /* { dg-error "invalid use of" } */ -#pragma acc declare present (v9) /* { dg-error "invalid use of" } */ +#pragma acc declare present (v2) /* { dg-error "invalid use of" } */ } diff --git a/gcc/testsuite/c-c++-common/goacc/default-4.c b/gcc/testsuite/c-c++-common/goacc/default-4.c index dfa79bbbe6e..867175d4847 100644 --- a/gcc/testsuite/c-c++-common/goacc/default-4.c +++ b/gcc/testsuite/c-c++-common/goacc/default-4.c @@ -8,7 +8,7 @@ void f1 () float f1_b[2]; #pragma acc data copyin (f1_a) copyout (f1_b) - /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f1_b \[^\\)\]+\\) map\\(force_to:f1_a" 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f1_b \[^\\)\]+\\) map\\(to:f1_a" 1 "gimple" } } */ { #pragma acc kernels /* { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f1_b \[^\\)\]+\\) map\\(tofrom:f1_a" 1 "gimple" } } */ @@ -29,7 +29,7 @@ void f2 () float f2_b[2]; #pragma acc data copyin (f2_a) copyout (f2_b) - /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f2_b \[^\\)\]+\\) map\\(force_to:f2_a" 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f2_b \[^\\)\]+\\) map\\(to:f2_a" 1 "gimple" } } */ { #pragma acc kernels default (none) /* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(none\\) map\\(tofrom:f2_b \[^\\)\]+\\) map\\(tofrom:f2_a" 1 "gimple" } } */ @@ -50,7 +50,7 @@ void f3 () float f3_b[2]; #pragma acc data copyin (f3_a) copyout (f3_b) - /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f3_b \[^\\)\]+\\) map\\(force_to:f3_a" 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f3_b \[^\\)\]+\\) map\\(to:f3_a" 1 "gimple" } } */ { #pragma acc kernels default (present) /* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/finalize-1.c b/gcc/testsuite/c-c++-common/goacc/finalize-1.c new file mode 100644 index 00000000000..94820290b94 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/finalize-1.c @@ -0,0 +1,28 @@ +/* Test valid usage and processing of the finalize clause. */ + +/* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */ + +extern int del_r; +extern float del_f[3]; +extern double cpo_r[8]; +extern long cpo_f; + +void f () +{ +#pragma acc exit data delete (del_r) +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_r\\);$" 1 "original" } } + { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */ + +#pragma acc exit data finalize delete (del_f) +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } } + { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } */ + +#pragma acc exit data copyout (cpo_r) +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } } + { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */ + +#pragma acc exit data copyout (cpo_f) finalize +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:cpo_f\\);$" 1 "original" } } + { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */ +} + diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c index d437c47779d..7576a6484f1 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c @@ -18,10 +18,12 @@ foo (void) } } +/* The xfails occur due to the OpenACC 2.5 data semantics. */ + /* { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" { xfail *-*-* } } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" { xfail *-*-* } } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" { xfail *-*-* } } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" { xfail *-*-* } } } */ /* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias.c index 25821ab2aea..e8ff018d512 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-alias.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias.c @@ -20,10 +20,12 @@ foo (void) } } +/* The xfails occur due to the OpenACC 2.5 data semantics. */ + /* { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" { xfail *-*-* } } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" { xfail *-*-* } } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" { xfail *-*-* } } } */ +/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" { xfail *-*-* } } } */ /* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/routine-5.c b/gcc/testsuite/c-c++-common/goacc/routine-5.c index b967a7447bd..b759db3292d 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-5.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-5.c @@ -4,11 +4,11 @@ struct PC { -#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */ }; void PC1( /* { dg-bogus "variable or field .PC1. declared void" "TODO" { xfail c++ } } */ -#pragma acc routine +#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" "" { target c } .-1 } { dg-error ".#pragma. is not allowed here" "" { target c++ } .-2 } */ ) /* { dg-bogus "expected declaration specifiers or .\\.\\.\\.. before .\\). token" "TODO" { xfail c } } */ @@ -18,26 +18,26 @@ void PC1( /* { dg-bogus "variable or field .PC1. declared void" "TODO" { xfail c void PC2() { if (0) -#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */ ; } void PC3() { -#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */ } /* "( name )" syntax. */ #pragma acc routine ( /* { dg-error "expected (function name|unqualified-id) before end of line" } */ -#pragma acc routine () /* { dg-error "expected (function name|unqualified-id) before .\\). token" } */ -#pragma acc routine (+) /* { dg-error "expected (function name|unqualified-id) before .\\+. token" } */ -#pragma acc routine (?) /* { dg-error "expected (function name|unqualified-id) before .\\?. token" } */ -#pragma acc routine (:) /* { dg-error "expected (function name|unqualified-id) before .:. token" } */ -#pragma acc routine (4) /* { dg-error "expected (function name|unqualified-id) before numeric constant" } */ +#pragma acc routine () seq /* { dg-error "expected (function name|unqualified-id) before .\\). token" } */ +#pragma acc routine (+) seq /* { dg-error "expected (function name|unqualified-id) before .\\+. token" } */ +#pragma acc routine (?) seq /* { dg-error "expected (function name|unqualified-id) before .\\?. token" } */ +#pragma acc routine (:) seq /* { dg-error "expected (function name|unqualified-id) before .:. token" } */ +#pragma acc routine (4) seq /* { dg-error "expected (function name|unqualified-id) before numeric constant" } */ #pragma acc routine ('4') /* { dg-error "expected (function name|unqualified-id) before .4." } */ -#pragma acc routine ("4") /* { dg-error "expected (function name|unqualified-id) before string constant" } */ +#pragma acc routine ("4") seq /* { dg-error "expected (function name|unqualified-id) before string constant" } */ extern void R1(void); extern void R2(void); #pragma acc routine (R1, R2, R3) worker /* { dg-error "expected .\\). before .,. token" } */ @@ -49,84 +49,84 @@ extern void R2(void); /* "#pragma acc routine" not immediately followed by (a single) function declaration or definition. */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ int a; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ void fn1 (void), fn1b (void); -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ int b, fn2 (void); -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ int b_, fn2_ (void), B_; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ int fn3 (void), b2; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ typedef struct c c; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ struct d {} d; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ void fn1_2 (void), fn1b_2 (void); -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ #pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ int b_2, fn2_2 (void); -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ int b_2_, fn2_2_ (void), B_2_; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */ int fn3_2 (void), b2_2; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ typedef struct c_2 c_2; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ struct d_2 {} d_2; -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq int fn4 (void); int fn5a (void); int fn5b (void); -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine (fn5a) -#pragma acc routine (fn5b) +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine (fn5a) seq +#pragma acc routine (fn5b) seq int fn5 (void); -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ -#pragma acc routine (fn6a) /* { dg-error ".fn6a. has not been declared" } */ -#pragma acc routine (fn6b) /* { dg-error ".fn6b. has not been declared" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine (fn6a) seq /* { dg-error ".fn6a. has not been declared" } */ +#pragma acc routine (fn6b) seq /* { dg-error ".fn6b. has not been declared" } */ int fn6 (void); #ifdef __cplusplus -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */ namespace f {} namespace g {} -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */ using namespace g; -#pragma acc routine (g) /* { dg-error ".g. does not refer to a function" "" { target c++ } } */ +#pragma acc routine (g) seq /* { dg-error ".g. does not refer to a function" "" { target c++ } } */ #endif /* __cplusplus */ -#pragma acc routine (a) /* { dg-error ".a. does not refer to a function" } */ +#pragma acc routine (a) seq /* { dg-error ".a. does not refer to a function" } */ -#pragma acc routine (c) /* { dg-error ".c. does not refer to a function" } */ +#pragma acc routine (c) seq /* { dg-error ".c. does not refer to a function" } */ /* Static assert. */ @@ -143,66 +143,24 @@ static_assert(0, ""); /* { dg-error "static assertion failed" "" { target c++11 #endif void f_static_assert(); /* Check that we already recognized "f_static_assert" as an OpenACC routine. */ -#pragma acc routine (f_static_assert) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*f_static_assert" "TODO" { xfail *-*-* } } */ +#pragma acc routine (f_static_assert) seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*f_static_assert" "TODO" { xfail *-*-* } } */ /* __extension__ usage. */ -#pragma acc routine +#pragma acc routine seq __extension__ extern void ex1(); #pragma acc routine (ex1) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*ex1" } */ -#pragma acc routine +#pragma acc routine seq __extension__ __extension__ __extension__ __extension__ __extension__ void ex2() { } #pragma acc routine (ex2) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*ex2" } */ -#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ +#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */ __extension__ int ex3; -#pragma acc routine (ex3) /* { dg-error ".ex3. does not refer to a function" } */ - - -/* "#pragma acc routine" already applied. */ - -extern void fungsi_1(); -#pragma acc routine(fungsi_1) gang -#pragma acc routine(fungsi_1) gang /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */ -#pragma acc routine(fungsi_1) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */ -#pragma acc routine(fungsi_1) vector /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */ - -#pragma acc routine seq -extern void fungsi_2(); -#pragma acc routine(fungsi_2) seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */ -#pragma acc routine(fungsi_2) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */ -#pragma acc routine(fungsi_2) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */ - -#pragma acc routine vector -extern void fungsi_3(); -#pragma acc routine vector /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_3." } */ -void fungsi_3() -{ -} - -extern void fungsi_4(); -#pragma acc routine (fungsi_4) worker -#pragma acc routine gang /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_4." } */ -void fungsi_4() -{ -} - -#pragma acc routine gang -void fungsi_5() -{ -} -#pragma acc routine (fungsi_5) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_5." } */ - -#pragma acc routine seq -void fungsi_6() -{ -} -#pragma acc routine seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_6." } */ -extern void fungsi_6(); +#pragma acc routine (ex3) seq /* { dg-error ".ex3. does not refer to a function" } */ /* "#pragma acc routine" must be applied before. */ @@ -214,11 +172,11 @@ void Foo () Bar (); } -#pragma acc routine (Bar) // { dg-error ".#pragma acc routine. must be applied before use" } +#pragma acc routine (Bar) seq // { dg-error ".#pragma acc routine. must be applied before use" } #pragma acc routine (Foo) gang // { dg-error ".#pragma acc routine. must be applied before definition" } -#pragma acc routine (Baz) // { dg-error "not been declared" } +#pragma acc routine (Baz) seq // { dg-error "not been declared" } /* OpenACC declare. */ @@ -227,7 +185,7 @@ int vb1; /* { dg-error "directive for use" } */ extern int vb2; /* { dg-error "directive for use" } */ static int vb3; /* { dg-error "directive for use" } */ -#pragma acc routine +#pragma acc routine seq int func1 (int a) { @@ -238,7 +196,7 @@ func1 (int a) return vb3; } -#pragma acc routine +#pragma acc routine seq int func2 (int a) { @@ -256,7 +214,7 @@ extern int vb6; /* { dg-error "clause used in" } */ static int vb7; /* { dg-error "clause used in" } */ #pragma acc declare link (vb7) -#pragma acc routine +#pragma acc routine seq int func3 (int a) { @@ -273,7 +231,7 @@ extern int vb9; static int vb10; #pragma acc declare create (vb10) -#pragma acc routine +#pragma acc routine seq int func4 (int a) { @@ -291,7 +249,7 @@ extern int vb12; extern int vb13; #pragma acc declare device_resident (vb13) -#pragma acc routine +#pragma acc routine seq int func5 (int a) { @@ -302,7 +260,7 @@ func5 (int a) return vb13; } -#pragma acc routine +#pragma acc routine seq int func6 (int a) { diff --git a/gcc/testsuite/c-c++-common/goacc/update-if_present-1.c b/gcc/testsuite/c-c++-common/goacc/update-if_present-1.c new file mode 100644 index 00000000000..c34a0e48065 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/update-if_present-1.c @@ -0,0 +1,28 @@ +/* Test valid usages of the if_present clause. */ + +/* { dg-additional-options "-fdump-tree-omplower" } */ + +void +t () +{ + int a, b, c[10]; + +#pragma acc update self(a) if_present +#pragma acc update device(b) async if_present +#pragma acc update host(c[1:3]) wait(4) if_present +#pragma acc update self(c) device(b) host (a) async(10) if (a == 5) if_present + +#pragma acc update self(a) +#pragma acc update device(b) async +#pragma acc update host(c[1:3]) wait(4) +#pragma acc update self(c) device(b) host (a) async(10) if (a == 5) +} + +/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present map.from:a .len: 4.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present async.-1. map.to:b .len: 4.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present wait.4. map.from:c.1. .len: 12.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present if.... async.10. map.from:a .len: 4.. map.to:b .len: 4.. map.from:c .len: 40.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "omp target oacc_update map.force_from:a .len: 4.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "omp target oacc_update async.-1. map.force_to:b .len: 4.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "omp target oacc_update wait.4. map.force_from:c.1. .len: 12.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "omp target oacc_update if.... async.10. map.force_from:a .len: 4.. map.force_to:b .len: 4.. map.force_from:c .len: 40.." 1 "omplower" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/update-if_present-2.c b/gcc/testsuite/c-c++-common/goacc/update-if_present-2.c new file mode 100644 index 00000000000..974f1b8c427 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/update-if_present-2.c @@ -0,0 +1,42 @@ +/* Test invalid usages of the if_present clause. */ + +#pragma acc routine gang if_present /* { dg-error "'if_present' is not valid" } */ +void +t1 () +{ + int a, b, c[10]; + +#pragma acc enter data copyin(a) if_present /* { dg-error "'if_present' is not valid" } */ +#pragma acc exit data copyout(a) if_present /* { dg-error "'if_present' is not valid" } */ + +#pragma acc data copy(a) if_present /* { dg-error "'if_present' is not valid" } */ + { + } + +#pragma acc declare create(c) if_present /* { dg-error "'if_present' is not valid" } */ + +#pragma acc init if_present +#pragma acc shutdown if_present +} + +void +t2 () +{ + int a, b, c[10]; + +#pragma acc update self(a) +#pragma acc parallel +#pragma acc loop if_present /* { dg-error "'if_present' is not valid" } */ + for (b = 1; b < 10; b++) + ; +#pragma acc end parallel + +#pragma acc kernels loop if_present /* { dg-error "'if_present' is not valid" } */ + for (b = 1; b < 10; b++) + ; + +#pragma acc parallel loop if_present /* { dg-error "'if_present' is not valid" } */ + for (b = 1; b < 10; b++) + ; +} + diff --git a/gcc/testsuite/g++.dg/goacc/template.C b/gcc/testsuite/g++.dg/goacc/template.C index 852f42f2b42..dae92b08987 100644 --- a/gcc/testsuite/g++.dg/goacc/template.C +++ b/gcc/testsuite/g++.dg/goacc/template.C @@ -1,4 +1,4 @@ -#pragma acc routine +#pragma acc routine seq template <typename T> T accDouble(int val) { @@ -31,7 +31,7 @@ oacc_parallel_copy (T a) #pragma acc parallel num_gangs (a) if (1) { -#pragma acc loop independent collapse (2) gang +#pragma acc loop independent collapse (2) for (int i = 0; i < a; i++) for (int j = 0; j < 5; j++) b = a; @@ -86,6 +86,8 @@ oacc_parallel_copy (T a) #pragma acc update self (b) #pragma acc update device (b) #pragma acc exit data delete (b) +#pragma acc exit data finalize copyout (b) +#pragma acc exit data delete (b) finalize return b; } @@ -133,6 +135,13 @@ oacc_kernels_copy (T a) b = a; } +#pragma acc update host (b) +#pragma acc update self (b) +#pragma acc update device (b) +#pragma acc exit data delete (b) +#pragma acc exit data finalize copyout (b) +#pragma acc exit data delete (b) finalize + return b; } diff --git a/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 index 42a447ad06b..956349204f4 100644 --- a/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 @@ -146,5 +146,5 @@ end subroutine test ! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. tile.2, 3" 2 "gimple" } } ! { dg-final { scan-tree-dump-times "acc loop private.i. independent" 2 "gimple" } } ! { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } } -! { dg-final { scan-tree-dump-times "omp target oacc_\[^ \]+ map.force_tofrom:y" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "omp target oacc_\[^ \]+ map.tofrom:y" 2 "gimple" } } ! { dg-final { scan-tree-dump-times "acc loop private.i. reduction..:y." 2 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/data-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/data-tree.f95 index 44efc8a670b..f16d62cce69 100644 --- a/gcc/testsuite/gfortran.dg/goacc/data-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/data-tree.f95 @@ -15,10 +15,10 @@ end program test ! { dg-final { scan-tree-dump-times "pragma acc data" 1 "original" } } ! { dg-final { scan-tree-dump-times "if" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_from:k\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_alloc:m\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/declare-2.f95 b/gcc/testsuite/gfortran.dg/goacc/declare-2.f95 index aa1704f77d0..7aa3dab4707 100644 --- a/gcc/testsuite/gfortran.dg/goacc/declare-2.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/declare-2.f95 @@ -11,11 +11,11 @@ subroutine asubr (b) !$acc declare copyout (b) ! { dg-error "Invalid clause in module" } !$acc declare present (b) ! { dg-error "Invalid clause in module" } !$acc declare present_or_copy (b) ! { dg-error "Invalid clause in module" } - !$acc declare present_or_copyin (b) ! { dg-error "Invalid clause in module" } + !$acc declare present_or_copyin (b) ! { dg-error "present on multiple" } !$acc declare present_or_copyout (b) ! { dg-error "Invalid clause in module" } - !$acc declare present_or_create (b) ! { dg-error "Invalid clause in module" } + !$acc declare present_or_create (b) ! { dg-error "present on multiple" } !$acc declare deviceptr (b) ! { dg-error "Invalid clause in module" } - !$acc declare create (b) copyin (b) ! { dg-error "present on multiple clauses" } + !$acc declare create (b) copyin (b) ! { dg-error "present on multiple" } end subroutine diff --git a/gcc/testsuite/gfortran.dg/goacc/default-4.f b/gcc/testsuite/gfortran.dg/goacc/default-4.f index 77291f43eff..30f411f70ab 100644 --- a/gcc/testsuite/gfortran.dg/goacc/default-4.f +++ b/gcc/testsuite/gfortran.dg/goacc/default-4.f @@ -8,7 +8,7 @@ REAL, DIMENSION (2) :: F1_B !$ACC DATA COPYIN (F1_A) COPYOUT (F1_B) -! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f1_a \[^\\)\]+\\) map\\(force_from:f1_b" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f1_a \[^\\)\]+\\) map\\(from:f1_b" 1 "gimple" } } !$ACC KERNELS ! { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f1_b \[^\\)\]+\\) map\\(tofrom:f1_a" 1 "gimple" } } F1_B(1) = F1_A; @@ -26,7 +26,7 @@ REAL, DIMENSION (2) :: F2_B !$ACC DATA COPYIN (F2_A) COPYOUT (F2_B) -! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f2_a \[^\\)\]+\\) map\\(force_from:f2_b" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f2_a \[^\\)\]+\\) map\\(from:f2_b" 1 "gimple" } } !$ACC KERNELS DEFAULT (NONE) ! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(none\\) map\\(tofrom:f2_b \[^\\)\]+\\) map\\(tofrom:f2_a" 1 "gimple" } } F2_B(1) = F2_A; @@ -44,7 +44,7 @@ REAL, DIMENSION (2) :: F3_B !$ACC DATA COPYIN (F3_A) COPYOUT (F3_B) -! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f3_a \[^\\)\]+\\) map\\(force_from:f3_b" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f3_a \[^\\)\]+\\) map\\(from:f3_b" 1 "gimple" } } !$ACC KERNELS DEFAULT (PRESENT) ! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } F3_B(1) = F3_A; diff --git a/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95 b/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95 index 8f1715e659d..805459c1bb0 100644 --- a/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95 @@ -84,5 +84,8 @@ contains !$acc exit data delete (tip) ! { dg-error "POINTER" } !$acc exit data delete (tia) ! { dg-error "ALLOCATABLE" } !$acc exit data copyout (i) delete (i) ! { dg-error "multiple clauses" } + !$acc exit data finalize + !$acc exit data finalize copyout (i) + !$acc exit data finalize delete (i) end subroutine foo end module test diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f new file mode 100644 index 00000000000..5c7a921a2e3 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f @@ -0,0 +1,27 @@ +! Test valid usage and processing of the finalize clause. + +! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } + + SUBROUTINE f + IMPLICIT NONE + INTEGER :: del_r + REAL, DIMENSION (3) :: del_f + DOUBLE PRECISION, DIMENSION (8) :: cpo_r + LOGICAL :: cpo_f + +!$ACC EXIT DATA DELETE (del_r) +! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_r\\);$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } + +!$ACC EXIT DATA FINALIZE DELETE (del_f) +! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } + +!$ACC EXIT DATA COPYOUT (cpo_r) +! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } + +!$ACC EXIT DATA COPYOUT (cpo_f) FINALIZE +! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_f\\) finalize;$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } + END SUBROUTINE f diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-alias-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-alias-2.f95 index 7e348dde2bd..6a9f241a596 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-alias-2.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-alias-2.f95 @@ -15,9 +15,11 @@ program main end program main +! The xfails occur in light of the new OpenACC data semantics. + ! { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } -! { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } -! { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } -! { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } -! { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } +! { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" { xfail *-*-* } } } +! { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" { xfail *-*-* } } } +! { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" { xfail *-*-* } } } +! { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" { xfail *-*-* } } } ! { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-alias.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-alias.f95 index 8d6ccb338b9..62f9a713991 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-alias.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-alias.f95 @@ -15,9 +15,11 @@ program main end program main +! The xfails occur in light of the new OpenACC data semantics. + ! { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } -! { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } -! { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } -! { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } -! { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } +! { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" { xfail *-*-* } } } +! { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" { xfail *-*-* } } } +! { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" { xfail *-*-* } } } +! { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" { xfail *-*-* } } } ! { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 index 7daca59020e..a70f1e737bd 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 @@ -21,10 +21,10 @@ end program test ! { dg-final { scan-tree-dump-times "num_workers" 1 "original" } } ! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_from:k\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_alloc:m\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 index 2fcaa400ee3..005193f30a7 100644 --- a/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 @@ -25,6 +25,8 @@ contains local_a (:) = 5 local_arg = 5 + !$acc update device(local_a) if_present + !$acc kernels loop & !$acc gang(num:local_arg) worker(local_arg) vector(local_arg) & !$acc wait async(local_arg) @@ -54,12 +56,16 @@ contains enddo enddo !$acc end kernels loop + + !$acc exit data copyout(local_a) delete(local_i) finalize end subroutine local subroutine nonlocal () nonlocal_a (:) = 5 nonlocal_arg = 5 + !$acc update device(nonlocal_a) if_present + !$acc kernels loop & !$acc gang(num:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) & !$acc wait async(nonlocal_arg) @@ -89,5 +95,7 @@ contains enddo enddo !$acc end kernels loop + + !$acc exit data copyout(nonlocal_a) delete(nonlocal_i) finalize end subroutine nonlocal end program main diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 index 5b2e01d4878..2697bb79e7f 100644 --- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 @@ -1,5 +1,4 @@ -! { dg-do compile } -! { dg-additional-options "-fdump-tree-original" } +! { dg-additional-options "-fdump-tree-original" } ! test for tree-dump-original and spaces-commas @@ -15,6 +14,7 @@ program test !$acc end parallel end program test + ! { dg-final { scan-tree-dump-times "pragma acc parallel" 1 "original" } } ! { dg-final { scan-tree-dump-times "if" 1 "original" } } @@ -24,10 +24,10 @@ end program test ! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } } ! { dg-final { scan-tree-dump-times "reduction\\(max:q\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_from:k\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "map\\(force_alloc:m\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90 b/gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90 index 6ff913ade8d..1d247ca238e 100644 --- a/gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90 @@ -38,9 +38,7 @@ program test !$acc end parallel end program test -! { dg-final { scan-tree-dump-times "map.tofrom:v1" 8 "gimple" } } -! { dg-final { scan-tree-dump-times "map.tofrom:v2" 8 "gimple" } } -! { dg-final { scan-tree-dump-times "map.force_tofrom:v1" 1 "gimple" } } -! { dg-final { scan-tree-dump-times "map.force_tofrom:v2" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "map.tofrom:v1" 9 "gimple" } } +! { dg-final { scan-tree-dump-times "map.tofrom:v2" 9 "gimple" } } ! { dg-final { scan-tree-dump-times "map.force_present:v1" 1 "gimple" } } ! { dg-final { scan-tree-dump-times "map.force_present:v2" 1 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/update-if_present-1.f90 b/gcc/testsuite/gfortran.dg/goacc/update-if_present-1.f90 new file mode 100644 index 00000000000..a183aae44c5 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/update-if_present-1.f90 @@ -0,0 +1,27 @@ +! Test valid usages of the if_present clause. + +! { dg-additional-options "-fdump-tree-omplower" } + +subroutine t + implicit none + integer a, b, c(10) + real, allocatable :: x, y, z(:) + + a = 5 + b = 10 + c(:) = -1 + + allocate (x, y, z(100)) + + !$acc update self(a) if_present + !$acc update device(b) if_present async + !$acc update host(c(1:3)) wait(4) if_present + !$acc update self(c) device(a) host(b) if_present async(10) if(a == 10) + + !$acc update self(x) if_present + !$acc update device(y) if_present async + !$acc update host(z(1:3)) wait(3) if_present + !$acc update self(z) device(y) host(x) if_present async(4) if(a == 1) +end subroutine t + +! { dg-final { scan-tree-dump-times " if_present" 8 "omplower" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/update-if_present-2.f90 b/gcc/testsuite/gfortran.dg/goacc/update-if_present-2.f90 new file mode 100644 index 00000000000..e73c2dc0875 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/update-if_present-2.f90 @@ -0,0 +1,52 @@ +! Test invalid usages of the if_present clause. + +subroutine t1 + implicit none + !$acc routine gang if_present ! { dg-error "Unclassifiable OpenACC directive" } + integer a, b, c(10) + real, allocatable :: x, y, z(:) + + a = 5 + b = 10 + c(:) = -1 + + allocate (x, y, z(100)) + + !$acc enter data copyin(a) if_present ! { dg-error "Unclassifiable OpenACC directive" } + !$acc exit data copyout(a) if_present ! { dg-error "Unclassifiable OpenACC directive" } + + !$acc data copy(a) if_present ! { dg-error "Unclassifiable OpenACC directive" } + !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" } + + !$acc declare link(a) if_present ! { dg-error "Unexpected junk after" } + + !$acc init if_present ! { dg-error "Unclassifiable OpenACC directive" } + !$acc shutdown if_present ! { dg-error "Unclassifiable OpenACC directive" } + + !$acc update self(a) device_type(nvidia) device(b) if_present ! { dg-error "Unclassifiable OpenACC directive" } +end subroutine t1 + +subroutine t2 + implicit none + integer a, b, c(10) + + a = 5 + b = 10 + c(:) = -1 + + !$acc parallel + !$acc loop if_present ! { dg-error "Unclassifiable OpenACC directive" } + do b = 1, 10 + end do + !$acc end parallel + + !$acc kernels loop if_present ! { dg-error "Unclassifiable OpenACC directive" } + do b = 1, 10 + end do + !$acc end kernels loop ! { dg-error "Unexpected ..ACC END KERNELS LOOP statement" } + + !$acc parallel loop if_present ! { dg-error "Unclassifiable OpenACC directive" } + do b = 1, 10 + end do + !$acc end parallel loop ! { dg-error "Unexpected ..ACC END PARALLEL LOOP statement" } +end subroutine t2 -- 2.17.1