Hi! Ping.
On Thu, 11 May 2017 14:26:51 +0200, I wrote: > [...] support the num_gangs, num_workers, vector_length > clauses for the OpenACC 2.5 kernels construct. OK for trunk? > > commit a689c52cde71960bc08ae30c3f88980f66fdd0b8 > Author: Thomas Schwinge <tho...@codesourcery.com> > Date: Thu May 11 13:43:28 2017 +0200 > > OpenACC 2.5 kernels construct: num_gangs, num_workers, vector_length > clauses > > gcc/c/ > * c-parser.c (OACC_KERNELS_CLAUSE_MASK): Add > "PRAGMA_OACC_CLAUSE_NUM_GANGS", "PRAGMA_OACC_CLAUSE_NUM_WORKERS", > "VECTOR_LENGTH". > gcc/cp/ > * parser.c (OACC_KERNELS_CLAUSE_MASK): Add > "PRAGMA_OACC_CLAUSE_NUM_GANGS", "PRAGMA_OACC_CLAUSE_NUM_WORKERS", > "VECTOR_LENGTH". > gcc/fortran/ > * openmp.c (OACC_KERNELS_CLAUSES): Add "OMP_CLAUSE_NUM_GANGS", > "OMP_CLAUSE_NUM_WORKERS", "OMP_CLAUSE_VECTOR_LENGTH". > gcc/ > * omp-offload.c (execute_oacc_device_lower): Remove the > parallelism dimensions function attributes for unparallelized > OpenACC kernels constructs. > gcc/testsuite/ > * c-c++-common/goacc/parallel-dims-1.c: Update. > * c-c++-common/goacc/parallel-dims-2.c: Likewise. > * c-c++-common/goacc/routine-1.c: Likewise. > * c-c++-common/goacc/uninit-dim-clause.c: Likewise. > * g++.dg/goacc/template.C: Likewise. > * gfortran.dg/goacc/kernels-tree.f95: Likewise. > * gfortran.dg/goacc/routine-3.f90: Likewise. > * gfortran.dg/goacc/sie.f95: Likewise. > * gfortran.dg/goacc/uninit-dim-clause.f95: Likewise. > libgomp/ > * testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c: Update. > * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise. > * testsuite/libgomp.oacc-fortran/kernels-loop-2.f95: Likewise. > --- > gcc/c/c-parser.c | 3 + > gcc/cp/parser.c | 3 + > gcc/fortran/openmp.c | 3 +- > gcc/omp-offload.c | 9 ++ > gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c | 3 + > gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c | 152 > +++++++++++++++++++-- > gcc/testsuite/c-c++-common/goacc/routine-1.c | 7 + > .../c-c++-common/goacc/uninit-dim-clause.c | 20 ++- > gcc/testsuite/g++.dg/goacc/template.C | 4 + > gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 | 6 +- > gcc/testsuite/gfortran.dg/goacc/routine-3.f90 | 6 + > gcc/testsuite/gfortran.dg/goacc/sie.f95 | 86 +++++++++++- > .../gfortran.dg/goacc/uninit-dim-clause.f95 | 18 ++- > .../libgomp.oacc-c-c++-common/kernels-loop-2.c | 21 ++- > .../libgomp.oacc-c-c++-common/parallel-dims.c | 35 +++++ > .../libgomp.oacc-fortran/kernels-loop-2.f95 | 13 +- > 16 files changed, 358 insertions(+), 31 deletions(-) > > diff --git gcc/c/c-parser.c gcc/c/c-parser.c > index 90d2d17..c0d733c 100644 > --- gcc/c/c-parser.c > +++ gcc/c/c-parser.c > @@ -13978,11 +13978,14 @@ c_parser_oacc_loop (location_t loc, c_parser > *parser, char *p_name, > | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \ > | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ > | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ > + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \ > + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \ > | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ > | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ > | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ > | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ > | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ > + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ > | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) > > #define OACC_PARALLEL_CLAUSE_MASK \ > diff --git gcc/cp/parser.c gcc/cp/parser.c > index 17d2679..0578e81 100644 > --- gcc/cp/parser.c > +++ gcc/cp/parser.c > @@ -36438,11 +36438,14 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token > *pragma_tok, char *p_name, > | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \ > | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ > | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ > + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \ > + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \ > | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ > | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ > | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ > | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ > | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ > + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ > | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) > > #define OACC_PARALLEL_CLAUSE_MASK \ > diff --git gcc/fortran/openmp.c gcc/fortran/openmp.c > index 89eecfa..7b18a1d 100644 > --- gcc/fortran/openmp.c > +++ gcc/fortran/openmp.c > @@ -1926,7 +1926,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const > omp_mask mask, > | OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_DEVICEPTR | > OMP_CLAUSE_PRIVATE \ > | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT) > #define OACC_KERNELS_CLAUSES \ > - (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_DEVICEPTR > \ > + (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS > \ > + | OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | > OMP_CLAUSE_DEVICEPTR \ > | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT > \ > | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY > \ > | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \ > diff --git gcc/omp-offload.c gcc/omp-offload.c > index e954ee9..fab26d0 100644 > --- gcc/omp-offload.c > +++ gcc/omp-offload.c > @@ -1452,6 +1452,15 @@ execute_oacc_device_lower () > = (lookup_attribute ("oacc kernels parallelized", > DECL_ATTRIBUTES (current_function_decl)) != NULL); > > + /* Unparallelized OpenACC kernels constructs must get launched as 1 x 1 x 1 > + kernels, so remove the parallelism dimensions function attributes > + potentially set earlier on. */ > + if (is_oacc_kernels && !is_oacc_kernels_parallelized) > + { > + oacc_set_fn_attrib (current_function_decl, NULL, NULL); > + attrs = oacc_get_fn_attrib (current_function_decl); > + } > + > /* Discover, partition and process the loops. */ > oacc_loop *loops = oacc_loop_discovery (); > int fn_level = oacc_fn_attrib_level (attrs); > diff --git gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c > gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c > index a85d3d3..57f682f 100644 > --- gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c > +++ gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c > @@ -3,6 +3,9 @@ > > void f(int i) > { > +#pragma acc kernels num_gangs(i) num_workers(i) vector_length(i) > + ; > + > #pragma acc parallel num_gangs(i) num_workers(i) vector_length(i) > ; > } > diff --git gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c > gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c > index 30a3d17..acfbe7f 100644 > --- gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c > +++ gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c > @@ -1,18 +1,15 @@ > /* Invalid use of OpenACC parallelism dimensions clauses: num_gangs, > num_workers, vector_length. */ > > -void acc_kernels(int i) > +void f(int i, float f) > { > -#pragma acc kernels num_gangs(i) /* { dg-error "'num_gangs' is not valid for > '#pragma acc kernels'" } */ > +#pragma acc kernels num_gangs /* { dg-error "expected '\\(' before end of > line" } */ > ; > -#pragma acc kernels num_workers(i) /* { dg-error "'num_workers' is not valid > for '#pragma acc kernels'" } */ > +#pragma acc kernels num_workers /* { dg-error "expected '\\(' before end of > line" } */ > ; > -#pragma acc kernels vector_length(i) /* { dg-error "'vector_length' is not > valid for '#pragma acc kernels'" } */ > +#pragma acc kernels vector_length /* { dg-error "expected '\\(' before end > of line" } */ > ; > -} > > -void acc_parallel(int i, float f) > -{ > #pragma acc parallel num_gangs /* { dg-error "expected '\\(' before end of > line" } */ > ; > #pragma acc parallel num_workers /* { dg-error "expected '\\(' before end of > line" } */ > @@ -20,6 +17,14 @@ void acc_parallel(int i, float f) > #pragma acc parallel vector_length /* { dg-error "expected '\\(' before end > of line" } */ > ; > > + > +#pragma acc kernels num_gangs( /* { dg-error "expected (primary-|)expression > before end of line" } */ > + ; > +#pragma acc kernels num_workers( /* { dg-error "expected > (primary-|)expression before end of line" } */ > + ; > +#pragma acc kernels vector_length( /* { dg-error "expected > (primary-|)expression before end of line" } */ > + ; > + > #pragma acc parallel num_gangs( /* { dg-error "expected > (primary-|)expression before end of line" } */ > ; > #pragma acc parallel num_workers( /* { dg-error "expected > (primary-|)expression before end of line" } */ > @@ -27,6 +32,14 @@ void acc_parallel(int i, float f) > #pragma acc parallel vector_length( /* { dg-error "expected > (primary-|)expression before end of line" } */ > ; > > + > +#pragma acc kernels num_gangs() /* { dg-error "expected > (primary-|)expression before '\\)' token" } */ > + ; > +#pragma acc kernels num_workers() /* { dg-error "expected > (primary-|)expression before '\\)' token" } */ > + ; > +#pragma acc kernels vector_length() /* { dg-error "expected > (primary-|)expression before '\\)' token" } */ > + ; > + > #pragma acc parallel num_gangs() /* { dg-error "expected > (primary-|)expression before '\\)' token" } */ > ; > #pragma acc parallel num_workers() /* { dg-error "expected > (primary-|)expression before '\\)' token" } */ > @@ -34,6 +47,14 @@ void acc_parallel(int i, float f) > #pragma acc parallel vector_length() /* { dg-error "expected > (primary-|)expression before '\\)' token" } */ > ; > > + > +#pragma acc kernels num_gangs(1 /* { dg-error "expected '\\)' before end of > line" } */ > + ; > +#pragma acc kernels num_workers(1 /* { dg-error "expected '\\)' before end > of line" } */ > + ; > +#pragma acc kernels vector_length(1 /* { dg-error "expected '\\)' before end > of line" } */ > + ; > + > #pragma acc parallel num_gangs(1 /* { dg-error "expected '\\)' before end of > line" } */ > ; > #pragma acc parallel num_workers(1 /* { dg-error "expected '\\)' before end > of line" } */ > @@ -41,6 +62,14 @@ void acc_parallel(int i, float f) > #pragma acc parallel vector_length(1 /* { dg-error "expected '\\)' before > end of line" } */ > ; > > + > +#pragma acc kernels num_gangs(i /* { dg-error "expected '\\)' before end of > line" } */ > + ; > +#pragma acc kernels num_workers(i /* { dg-error "expected '\\)' before end > of line" } */ > + ; > +#pragma acc kernels vector_length(i /* { dg-error "expected '\\)' before end > of line" } */ > + ; > + > #pragma acc parallel num_gangs(i /* { dg-error "expected '\\)' before end of > line" } */ > ; > #pragma acc parallel num_workers(i /* { dg-error "expected '\\)' before end > of line" } */ > @@ -48,6 +77,14 @@ void acc_parallel(int i, float f) > #pragma acc parallel vector_length(i /* { dg-error "expected '\\)' before > end of line" } */ > ; > > + > +#pragma acc kernels num_gangs(1 i /* { dg-error "expected '\\)' before 'i'" > } */ > + ; > +#pragma acc kernels num_workers(1 i /* { dg-error "expected '\\)' before > 'i'" } */ > + ; > +#pragma acc kernels vector_length(1 i /* { dg-error "expected '\\)' before > 'i'" } */ > + ; > + > #pragma acc parallel num_gangs(1 i /* { dg-error "expected '\\)' before 'i'" > } */ > ; > #pragma acc parallel num_workers(1 i /* { dg-error "expected '\\)' before > 'i'" } */ > @@ -55,6 +92,14 @@ void acc_parallel(int i, float f) > #pragma acc parallel vector_length(1 i /* { dg-error "expected '\\)' before > 'i'" } */ > ; > > + > +#pragma acc kernels num_gangs(1 i) /* { dg-error "expected '\\)' before 'i'" > } */ > + ; > +#pragma acc kernels num_workers(1 i) /* { dg-error "expected '\\)' before > 'i'" } */ > + ; > +#pragma acc kernels vector_length(1 i) /* { dg-error "expected '\\)' before > 'i'" } */ > + ; > + > #pragma acc parallel num_gangs(1 i) /* { dg-error "expected '\\)' before > 'i'" } */ > ; > #pragma acc parallel num_workers(1 i) /* { dg-error "expected '\\)' before > 'i'" } */ > @@ -62,6 +107,17 @@ void acc_parallel(int i, float f) > #pragma acc parallel vector_length(1 i) /* { dg-error "expected '\\)' before > 'i'" } */ > ; > > + > +#pragma acc kernels num_gangs(1, i /* { dg-error "expected '\\)' before ',' > token" "TODO" { xfail c } } */ > + /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } > */ > + ; > +#pragma acc kernels num_workers(1, i /* { dg-error "expected '\\)' before > ',' token" "TODO" { xfail c } } */ > + /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } > */ > + ; > +#pragma acc kernels vector_length(1, i /* { dg-error "expected '\\)' before > ',' token" "TODO" { xfail c } } */ > + /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } > */ > + ; > + > #pragma acc parallel num_gangs(1, i /* { dg-error "expected '\\)' before ',' > token" "TODO" { xfail c } } */ > /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } > */ > ; > @@ -72,6 +128,14 @@ void acc_parallel(int i, float f) > /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } > */ > ; > > + > +#pragma acc kernels num_gangs(1, i) /* { dg-error "expected '\\)' before ',' > token" "TODO" { xfail c } } */ > + ; > +#pragma acc kernels num_workers(1, i) /* { dg-error "expected '\\)' before > ',' token" "TODO" { xfail c } } */ > + ; > +#pragma acc kernels vector_length(1, i) /* { dg-error "expected '\\)' before > ',' token" "TODO" { xfail c } } */ > + ; > + > #pragma acc parallel num_gangs(1, i) /* { dg-error "expected '\\)' before > ',' token" "TODO" { xfail c } } */ > ; > #pragma acc parallel num_workers(1, i) /* { dg-error "expected '\\)' before > ',' token" "TODO" { xfail c } } */ > @@ -79,11 +143,27 @@ void acc_parallel(int i, float f) > #pragma acc parallel vector_length(1, i) /* { dg-error "expected '\\)' > before ',' token" "TODO" { xfail c } } */ > ; > > -#pragma acc parallel num_gangs(num_gangs) /* { dg-error "'num_gangs' (un|was > not )declared" } */ > + > +#pragma acc kernels num_gangs(num_gangs_k) /* { dg-error "'num_gangs_k' > (un|was not )declared" } */ > ; > -#pragma acc parallel num_workers(num_workers) /* { dg-error "'num_workers' > (un|was not )declared" } */ > +#pragma acc kernels num_workers(num_workers_k) /* { dg-error > "'num_workers_k' (un|was not )declared" } */ > ; > -#pragma acc parallel vector_length(vector_length) /* { dg-error > "'vector_length' (un|was not )declared" } */ > +#pragma acc kernels vector_length(vector_length_k) /* { dg-error > "'vector_length_k' (un|was not )declared" } */ > + ; > + > +#pragma acc parallel num_gangs(num_gangs_p) /* { dg-error "'num_gangs_p' > (un|was not )declared" } */ > + ; > +#pragma acc parallel num_workers(num_workers_p) /* { dg-error > "'num_workers_p' (un|was not )declared" } */ > + ; > +#pragma acc parallel vector_length(vector_length_p) /* { dg-error > "'vector_length_p' (un|was not )declared" } */ > + ; > + > + > +#pragma acc kernels num_gangs(f) /* { dg-error "'num_gangs' expression must > be integral" } */ > + ; > +#pragma acc kernels num_workers(f) /* { dg-error "'num_workers' expression > must be integral" } */ > + ; > +#pragma acc kernels vector_length(f) /* { dg-error "'vector_length' > expression must be integral" } */ > ; > > #pragma acc parallel num_gangs(f) /* { dg-error "'num_gangs' expression must > be integral" } */ > @@ -93,6 +173,14 @@ void acc_parallel(int i, float f) > #pragma acc parallel vector_length(f) /* { dg-error "'vector_length' > expression must be integral" } */ > ; > > + > +#pragma acc kernels num_gangs((float) 1) /* { dg-error "'num_gangs' > expression must be integral" } */ > + ; > +#pragma acc kernels num_workers((float) 1) /* { dg-error "'num_workers' > expression must be integral" } */ > + ; > +#pragma acc kernels vector_length((float) 1) /* { dg-error "'vector_length' > expression must be integral" } */ > + ; > + > #pragma acc parallel num_gangs((float) 1) /* { dg-error "'num_gangs' > expression must be integral" } */ > ; > #pragma acc parallel num_workers((float) 1) /* { dg-error "'num_workers' > expression must be integral" } */ > @@ -100,6 +188,14 @@ void acc_parallel(int i, float f) > #pragma acc parallel vector_length((float) 1) /* { dg-error "'vector_length' > expression must be integral" } */ > ; > > + > +#pragma acc kernels num_gangs(0) /* { dg-warning "'num_gangs' value must be > positive" } */ > + ; > +#pragma acc kernels num_workers(0) /* { dg-warning "'num_workers' value must > be positive" } */ > + ; > +#pragma acc kernels vector_length(0) /* { dg-warning "'vector_length' value > must be positive" } */ > + ; > + > #pragma acc parallel num_gangs(0) /* { dg-warning "'num_gangs' value must be > positive" } */ > ; > #pragma acc parallel num_workers(0) /* { dg-warning "'num_workers' value > must be positive" } */ > @@ -107,6 +203,14 @@ void acc_parallel(int i, float f) > #pragma acc parallel vector_length(0) /* { dg-warning "'vector_length' value > must be positive" } */ > ; > > + > +#pragma acc kernels num_gangs((int) -1.2) /* { dg-warning "'num_gangs' value > must be positive" } */ > + ; > +#pragma acc kernels num_workers((int) -1.2) /* { dg-warning "'num_workers' > value must be positive" } */ > + ; > +#pragma acc kernels vector_length((int) -1.2) /* { dg-warning > "'vector_length' value must be positive" } */ > + ; > + > #pragma acc parallel num_gangs((int) -1.2) /* { dg-warning "'num_gangs' > value must be positive" } */ > ; > #pragma acc parallel num_workers((int) -1.2) /* { dg-warning "'num_workers' > value must be positive" } */ > @@ -114,7 +218,8 @@ void acc_parallel(int i, float f) > #pragma acc parallel vector_length((int) -1.2) /* { dg-warning > "'vector_length' value must be positive" } */ > ; > > -#pragma acc parallel \ > + > +#pragma acc kernels \ > num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c } > } */ \ > num_workers(1) /* { dg-error "too many 'num_workers' clauses" "" { target > c } } */ \ > vector_length(1) /* { dg-error "too many 'vector_length' clauses" "" { > target c } } */ \ > @@ -123,12 +228,31 @@ void acc_parallel(int i, float f) > num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c++ > } } */ > ; > > -#pragma acc parallel \ > +#pragma acc parallel \ > + num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c } > } */ \ > + num_workers(1) /* { dg-error "too many 'num_workers' clauses" "" { target > c } } */ \ > + vector_length(1) /* { dg-error "too many 'vector_length' clauses" "" { > target c } } */ \ > + num_workers(1) /* { dg-error "too many 'num_workers' clauses" "" { target > c++ } } */ \ > + vector_length(1) /* { dg-error "too many 'vector_length' clauses" "" { > target c++ } } */ \ > + num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c++ > } } */ > + ; > + > + > +#pragma acc kernels \ > + num_gangs(-1) /* { dg-warning "'num_gangs' value must be positive" } */ \ > + num_workers() /* { dg-error "expected (primary-|)expression before '\\)' > token" } */ \ > + vector_length(abc_k) /* { dg-error "'abc_k' (un|was not )declared" } */ \ > + num_workers(0.5) /* { dg-error "'num_workers' expression must be integral" > } */ \ > + vector_length(&f) /* { dg-error "'vector_length' expression must be > integral" } */ \ > + num_gangs( /* { dg-error "expected (primary-|)expression before end of > line" "TODO" { xfail c } } */ > + ; > + > +#pragma acc parallel \ > num_gangs(-1) /* { dg-warning "'num_gangs' value must be positive" } */ \ > num_workers() /* { dg-error "expected (primary-|)expression before '\\)' > token" } */ \ > - vector_length(abc) /* { dg-error "'abc' (un|was not )declared" } */ \ > + vector_length(abc_p) /* { dg-error "'abc_p' (un|was not )declared" } */ \ > num_workers(0.5) /* { dg-error "'num_workers' expression must be integral" > } */ \ > - vector_length(&acc_parallel) /* { dg-error "'vector_length' expression > must be integral" } */ \ > + vector_length(&f) /* { dg-error "'vector_length' expression must be > integral" } */ \ > num_gangs( /* { dg-error "expected (primary-|)expression before end of > line" "TODO" { xfail c } } */ > ; > } > diff --git gcc/testsuite/c-c++-common/goacc/routine-1.c > gcc/testsuite/c-c++-common/goacc/routine-1.c > index a5e0d69..a756922 100644 > --- gcc/testsuite/c-c++-common/goacc/routine-1.c > +++ gcc/testsuite/c-c++-common/goacc/routine-1.c > @@ -21,6 +21,13 @@ void seq (void) > > int main () > { > +#pragma acc kernels num_gangs (32) num_workers (32) vector_length (32) > + { > + gang (); > + worker (); > + vector (); > + seq (); > + } > > #pragma acc parallel num_gangs (32) num_workers (32) vector_length (32) > { > diff --git gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c > gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c > index 0a006e3..9f11196 100644 > --- gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c > +++ gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c > @@ -1,10 +1,6 @@ > -/* { dg-do compile } */ > /* { dg-additional-options "-Wuninitialized" } */ > > -#include <stdbool.h> > - > -int > -main (void) > +void acc_parallel() > { > int i, j, k; > > @@ -17,3 +13,17 @@ main (void) > #pragma acc parallel vector_length(k) /* { dg-warning "is used > uninitialized in this function" } */ > ; > } > + > +void acc_kernels() > +{ > + int i, j, k; > + > + #pragma acc kernels num_gangs(i) /* { dg-warning "is used uninitialized in > this function" } */ > + ; > + > + #pragma acc kernels num_workers(j) /* { dg-warning "is used uninitialized > in this function" } */ > + ; > + > + #pragma acc kernels vector_length(k) /* { dg-warning "is used > uninitialized in this function" } */ > + ; > +} > diff --git gcc/testsuite/g++.dg/goacc/template.C > gcc/testsuite/g++.dg/goacc/template.C > index 74f40d8..852f42f 100644 > --- gcc/testsuite/g++.dg/goacc/template.C > +++ gcc/testsuite/g++.dg/goacc/template.C > @@ -100,6 +100,10 @@ oacc_kernels_copy (T a) > float y = 3; > double z = 4; > > +#pragma acc kernels num_gangs (a) num_workers (a) vector_length (a) default > (none) copyout (b) copyin (a) > + for (int i = 0; i < 1; i++) > + b = a; > + > #pragma acc kernels copy (w, x, y, z) > { > w = accDouble<char>(w); > diff --git gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 > gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 > index 4ec66de..7daca59 100644 > --- gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 > +++ gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 > @@ -6,7 +6,8 @@ program test > integer :: q, i, j, k, m, n, o, p, r, s, t, u, v, w > logical :: l = .true. > > - !$acc kernels if(l) async copy(i), copyin(j), copyout(k), create(m) & > + !$acc kernels if(l) async num_gangs(i) num_workers(i) vector_length(i) & > + !$acc copy(i), copyin(j), copyout(k), create(m) & > !$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) & > !$acc deviceptr(u) > !$acc end kernels > @@ -16,6 +17,9 @@ end program test > > ! { dg-final { scan-tree-dump-times "if" 1 "original" } } > ! { dg-final { scan-tree-dump-times "async" 1 "original" } } > +! { dg-final { scan-tree-dump-times "num_gangs" 1 "original" } } > +! { 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" } } > diff --git gcc/testsuite/gfortran.dg/goacc/routine-3.f90 > gcc/testsuite/gfortran.dg/goacc/routine-3.f90 > index ca9b928..6773f62 100644 > --- gcc/testsuite/gfortran.dg/goacc/routine-3.f90 > +++ gcc/testsuite/gfortran.dg/goacc/routine-3.f90 > @@ -4,6 +4,12 @@ CONTAINS > INTEGER :: i > REAL(KIND=8), ALLOCATABLE :: un(:), ua(:) > > + !$acc kernels num_gangs(2) num_workers(4) vector_length(32) > + DO jj = 1, 100 > + un(i) = ua(i) > + END DO > + !$acc end kernels > + > !$acc parallel num_gangs(2) num_workers(4) vector_length(32) > DO jj = 1, 100 > un(i) = ua(i) > diff --git gcc/testsuite/gfortran.dg/goacc/sie.f95 > gcc/testsuite/gfortran.dg/goacc/sie.f95 > index 2d66026..abfe28b 100644 > --- gcc/testsuite/gfortran.dg/goacc/sie.f95 > +++ gcc/testsuite/gfortran.dg/goacc/sie.f95 > @@ -95,6 +95,34 @@ program test > !$acc parallel num_gangs("1") ! { dg-error "scalar INTEGER expression" } > !$acc end parallel > > + !$acc kernels num_gangs ! { dg-error "Unclassifiable OpenACC directive" } > + > + !$acc kernels num_gangs(3) > + !$acc end kernels > + > + !$acc kernels num_gangs(i) > + !$acc end kernels > + > + !$acc kernels num_gangs(i+1) > + !$acc end kernels > + > + !$acc kernels num_gangs(-1) ! { dg-warning "must be positive" } > + !$acc end kernels > + > + !$acc kernels num_gangs(0) ! { dg-warning "must be positive" } > + !$acc end kernels > + > + !$acc kernels num_gangs() ! { dg-error "Invalid character in name" } > + > + !$acc kernels num_gangs(1.5) ! { dg-error "scalar INTEGER expression" } > + !$acc end kernels > + > + !$acc kernels num_gangs(.true.) ! { dg-error "scalar INTEGER expression" } > + !$acc end kernels > + > + !$acc kernels num_gangs("1") ! { dg-error "scalar INTEGER expression" } > + !$acc end kernels > + > > !$acc parallel num_workers ! { dg-error "Unclassifiable OpenACC directive" > } > > @@ -124,6 +152,34 @@ program test > !$acc parallel num_workers("1") ! { dg-error "scalar INTEGER expression" } > !$acc end parallel > > + !$acc kernels num_workers ! { dg-error "Unclassifiable OpenACC directive" } > + > + !$acc kernels num_workers(3) > + !$acc end kernels > + > + !$acc kernels num_workers(i) > + !$acc end kernels > + > + !$acc kernels num_workers(i+1) > + !$acc end kernels > + > + !$acc kernels num_workers(-1) ! { dg-warning "must be positive" } > + !$acc end kernels > + > + !$acc kernels num_workers(0) ! { dg-warning "must be positive" } > + !$acc end kernels > + > + !$acc kernels num_workers() ! { dg-error "Invalid character in name" } > + > + !$acc kernels num_workers(1.5) ! { dg-error "scalar INTEGER expression" } > + !$acc end kernels > + > + !$acc kernels num_workers(.true.) ! { dg-error "scalar INTEGER expression" > } > + !$acc end kernels > + > + !$acc kernels num_workers("1") ! { dg-error "scalar INTEGER expression" } > + !$acc end kernels > + > > !$acc parallel vector_length ! { dg-error "Unclassifiable OpenACC > directive" } > > @@ -153,6 +209,34 @@ program test > !$acc parallel vector_length("1") ! { dg-error "scalar INTEGER expression" > } > !$acc end parallel > > + !$acc kernels vector_length ! { dg-error "Unclassifiable OpenACC > directive" } > + > + !$acc kernels vector_length(3) > + !$acc end kernels > + > + !$acc kernels vector_length(i) > + !$acc end kernels > + > + !$acc kernels vector_length(i+1) > + !$acc end kernels > + > + !$acc kernels vector_length(-1) ! { dg-warning "must be positive" } > + !$acc end kernels > + > + !$acc kernels vector_length(0) ! { dg-warning "must be positive" } > + !$acc end kernels > + > + !$acc kernels vector_length() ! { dg-error "Invalid character in name" } > + > + !$acc kernels vector_length(1.5) ! { dg-error "scalar INTEGER expression" } > + !$acc end kernels > + > + !$acc kernels vector_length(.true.) ! { dg-error "scalar INTEGER > expression" } > + !$acc end kernels > + > + !$acc kernels vector_length("1") ! { dg-error "scalar INTEGER expression" } > + !$acc end kernels > + > > !$acc loop gang > do i = 1,10 > @@ -249,4 +333,4 @@ program test > do i = 1,10 > enddo > > -end program test > \ No newline at end of file > +end program test > diff --git gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 > gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 > index b87d26f..5dea42b 100644 > --- gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 > +++ gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 > @@ -1,7 +1,6 @@ > -! { dg-do compile } > ! { dg-additional-options "-Wuninitialized" } > > -program test > +subroutine acc_parallel > implicit none > integer :: i, j, k > > @@ -13,5 +12,18 @@ program test > > !$acc parallel vector_length(k) ! { dg-warning "is used uninitialized in > this function" } > !$acc end parallel > +end subroutine acc_parallel > > -end program test > +subroutine acc_kernels > + implicit none > + integer :: i, j, k > + > + !$acc kernels num_gangs(i) ! { dg-warning "is used uninitialized in this > function" } > + !$acc end kernels > + > + !$acc kernels num_workers(j) ! { dg-warning "is used uninitialized in this > function" } > + !$acc end kernels > + > + !$acc kernels vector_length(k) ! { dg-warning "is used uninitialized in > this function" } > + !$acc end kernels > +end subroutine acc_kernels > diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c > libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c > index c7592d6..b840888 100644 > --- libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c > +++ libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c > @@ -14,27 +14,40 @@ main (void) > b = (unsigned int *__restrict)malloc (N * sizeof (unsigned int)); > c = (unsigned int *__restrict)malloc (N * sizeof (unsigned int)); > > + /* Parallelism dimensions: compiler/runtime decides. */ > #pragma acc kernels copyout (a[0:N]) > { > for (COUNTERTYPE i = 0; i < N; i++) > a[i] = i * 2; > } > > -#pragma acc kernels copyout (b[0:N]) > + /* Parallelism dimensions: variable. */ > +#pragma acc kernels copyout (b[0:N]) \ > + num_gangs (3 + a[3]) num_workers (5 + a[5]) vector_length (7 + a[7]) > + /* { dg-prune-output "using vector_length \\(32\\), ignoring runtime > setting" } */ > { > for (COUNTERTYPE i = 0; i < N; i++) > b[i] = i * 4; > } > > -#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) > + /* Parallelism dimensions: literal. */ > +#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) \ > + num_gangs (3) num_workers (5) vector_length (7) > + /* { dg-prune-output "using vector_length \\(32\\), ignoring 7" } */ > { > for (COUNTERTYPE ii = 0; ii < N; ii++) > c[ii] = a[ii] + b[ii]; > } > > for (COUNTERTYPE i = 0; i < N; i++) > - if (c[i] != a[i] + b[i]) > - abort (); > + { > + if (a[i] != i * 2) > + abort (); > + if (b[i] != i * 4) > + abort (); > + if (c[i] != a[i] + b[i]) > + abort (); > + } > > free (a); > free (b); > diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c > libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c > index d8af546..8308f7c 100644 > --- libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c > +++ libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c > @@ -520,5 +520,40 @@ int main () > } > > > + /* Unparallelized OpenACC kernels constructs must get launched as 1 x 1 x 1 > + kernels even when there are explicit num_gangs, num_workers, or > + vector_length clauses. */ > + { > + int gangs = 5; > +#define WORKERS 5 > +#define VECTORS 13 > + int gangs_min, gangs_max, workers_min, workers_max, vectors_min, > vectors_max; > + gangs_min = workers_min = vectors_min = INT_MAX; > + gangs_max = workers_max = vectors_max = INT_MIN; > +#pragma acc kernels \ > + num_gangs (gangs) \ > + num_workers (WORKERS) \ > + vector_length (VECTORS) > + { > + /* This is to make the OpenACC kernels construct unparallelizable. */ > + asm volatile ("" : : : "memory"); > + > +#pragma acc loop reduction (min: gangs_min, workers_min, vectors_min) > reduction (max: gangs_max, workers_max, vectors_max) > + for (int i = 100; i > -100; --i) > + { > + gangs_min = gangs_max = acc_gang (); > + workers_min = workers_max = acc_worker (); > + vectors_min = vectors_max = acc_vector (); > + } > + } > + if (gangs_min != 0 || gangs_max != 1 - 1 > + || workers_min != 0 || workers_max != 1 - 1 > + || vectors_min != 0 || vectors_max != 1 - 1) > + __builtin_abort (); > +#undef VECTORS > +#undef WORKERS > + } > + > + > return 0; > } > diff --git libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95 > libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95 > index 163e8d5..b88ca67 100644 > --- libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95 > +++ libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95 > @@ -6,25 +6,34 @@ program main > integer, dimension (0:n-1) :: a, b, c > integer :: i, ii > > + ! Parallelism dimensions: compiler/runtime decides. > !$acc kernels copyout (a(0:n-1)) > do i = 0, n - 1 > a(i) = i * 2 > end do > !$acc end kernels > > - !$acc kernels copyout (b(0:n-1)) > + ! Parallelism dimensions: variable. > + !$acc kernels copyout (b(0:n-1)) & > + !$acc num_gangs (3 + a(3)) num_workers (5 + a(5)) vector_length (7 + a(7)) > + ! { dg-prune-output "using vector_length \\(32\\), ignoring runtime > setting" } > do i = 0, n -1 > b(i) = i * 4 > end do > !$acc end kernels > > - !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) > + ! Parallelism dimensions: literal. > + !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) & > + !$acc num_gangs (3) num_workers (5) vector_length (7) > + ! { dg-prune-output "using vector_length \\(32\\), ignoring 7" } > do ii = 0, n - 1 > c(ii) = a(ii) + b(ii) > end do > !$acc end kernels > > do i = 0, n - 1 > + if (a(i) .ne. i * 2) call abort > + if (b(i) .ne. i * 4) call abort > if (c(i) .ne. a(i) + b(i)) call abort > end do Grüße Thomas