diff options
author | Thomas Schwinge <thomas@codesourcery.com> | 2017-05-23 17:47:32 +0200 |
---|---|---|
committer | Thomas Schwinge <tschwinge@gcc.gnu.org> | 2017-05-23 17:47:32 +0200 |
commit | fd71a9a24da0a52b9a752b4a7d4a6f71d96f5c29 (patch) | |
tree | 97c97af78e8ae6ff6ebb3d425ad65855616b9a43 /gcc | |
parent | 464d01188e46ccdb19b4783c5d28ef6bc72d32f9 (diff) | |
download | gcc-fd71a9a24da0a52b9a752b4a7d4a6f71d96f5c29.zip gcc-fd71a9a24da0a52b9a752b4a7d4a6f71d96f5c29.tar.gz gcc-fd71a9a24da0a52b9a752b4a7d4a6f71d96f5c29.tar.bz2 |
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.
From-SVN: r248370
Diffstat (limited to 'gcc')
-rw-r--r-- | gcc/ChangeLog | 6 | ||||
-rw-r--r-- | gcc/c/ChangeLog | 6 | ||||
-rw-r--r-- | gcc/c/c-parser.c | 3 | ||||
-rw-r--r-- | gcc/cp/ChangeLog | 6 | ||||
-rw-r--r-- | gcc/cp/parser.c | 3 | ||||
-rw-r--r-- | gcc/fortran/ChangeLog | 5 | ||||
-rw-r--r-- | gcc/fortran/openmp.c | 3 | ||||
-rw-r--r-- | gcc/omp-offload.c | 9 | ||||
-rw-r--r-- | gcc/testsuite/ChangeLog | 12 | ||||
-rw-r--r-- | gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c | 3 | ||||
-rw-r--r-- | gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c | 152 | ||||
-rw-r--r-- | gcc/testsuite/c-c++-common/goacc/routine-1.c | 7 | ||||
-rw-r--r-- | gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c | 20 | ||||
-rw-r--r-- | gcc/testsuite/g++.dg/goacc/template.C | 4 | ||||
-rw-r--r-- | gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 | 6 | ||||
-rw-r--r-- | gcc/testsuite/gfortran.dg/goacc/routine-3.f90 | 6 | ||||
-rw-r--r-- | gcc/testsuite/gfortran.dg/goacc/sie.f95 | 86 | ||||
-rw-r--r-- | gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 | 18 |
18 files changed, 330 insertions, 25 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index d2e846e..b38a31d 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,9 @@ +2017-05-23 Thomas Schwinge <thomas@codesourcery.com> + + * omp-offload.c (execute_oacc_device_lower): Remove the + parallelism dimensions function attributes for unparallelized + OpenACC kernels constructs. + 2017-05-23 Martin Liska <mliska@suse.cz> * cgraph.c (cgraph_node::get_create): Use symtab_node::dump_{asm_,}name diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index d768d93..cb04d4a 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,9 @@ +2017-05-23 Thomas Schwinge <thomas@codesourcery.com> + + * c-parser.c (OACC_KERNELS_CLAUSE_MASK): Add + "PRAGMA_OACC_CLAUSE_NUM_GANGS", "PRAGMA_OACC_CLAUSE_NUM_WORKERS", + "VECTOR_LENGTH". + 2017-05-23 Marek Polacek <polacek@redhat.com> * c-parser.c (c_parser_compound_statement_nostart): Remove redundant diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index f3bcbee..03c711b 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -13984,11 +13984,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 a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog index aaf39c2..bfe718f 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,3 +1,9 @@ +2017-05-23 Thomas Schwinge <thomas@codesourcery.com> + + * parser.c (OACC_KERNELS_CLAUSE_MASK): Add + "PRAGMA_OACC_CLAUSE_NUM_GANGS", "PRAGMA_OACC_CLAUSE_NUM_WORKERS", + "VECTOR_LENGTH". + 2017-05-23 Nathan Sidwell <nathan@acm.org> * cp-tree.h (OVL_P): New. diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 23d979c..b39e624 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -36432,11 +36432,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 a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog index 6977bd1..b3179e0 100644 --- a/gcc/fortran/ChangeLog +++ b/gcc/fortran/ChangeLog @@ -1,3 +1,8 @@ +2017-05-23 Thomas Schwinge <thomas@codesourcery.com> + + * openmp.c (OACC_KERNELS_CLAUSES): Add "OMP_CLAUSE_NUM_GANGS", + "OMP_CLAUSE_NUM_WORKERS", "OMP_CLAUSE_VECTOR_LENGTH". + 2017-05-22 Janus Weil <janus@gcc.gnu.org> PR fortran/80766 diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 80146e2..5a2b774 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -1932,7 +1932,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 a/gcc/omp-offload.c b/gcc/omp-offload.c index f02b4f8..54a4e90 100644 --- a/gcc/omp-offload.c +++ b/gcc/omp-offload.c @@ -1451,6 +1451,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 a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index 7b42abc..cef3ba6 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,15 @@ +2017-05-23 Thomas Schwinge <thomas@codesourcery.com> + + * 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. + 2017-05-23 Nathan Sidwell <nathan@acm.org> * g++.dg/lookup/using13.C: Adjust expected error. diff --git a/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c b/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c index a85d3d3..57f682f 100644 --- a/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c +++ b/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 a/gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c b/gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c index 30a3d17..acfbe7f 100644 --- a/gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c +++ b/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 kernels num_workers(num_workers_k) /* { dg-error "'num_workers_k' (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 parallel num_workers(num_workers) /* { dg-error "'num_workers' (un|was not )declared" } */ +#pragma acc kernels num_workers(f) /* { dg-error "'num_workers' expression must be integral" } */ ; -#pragma acc parallel vector_length(vector_length) /* { dg-error "'vector_length' (un|was not )declared" } */ +#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 a/gcc/testsuite/c-c++-common/goacc/routine-1.c b/gcc/testsuite/c-c++-common/goacc/routine-1.c index a5e0d69..a756922 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-1.c +++ b/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 a/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c b/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c index 0a006e3..9f11196 100644 --- a/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c +++ b/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 a/gcc/testsuite/g++.dg/goacc/template.C b/gcc/testsuite/g++.dg/goacc/template.C index 74f40d8..852f42f 100644 --- a/gcc/testsuite/g++.dg/goacc/template.C +++ b/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 a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 index 4ec66de..7daca59 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 +++ b/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 a/gcc/testsuite/gfortran.dg/goacc/routine-3.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-3.f90 index ca9b928..6773f62 100644 --- a/gcc/testsuite/gfortran.dg/goacc/routine-3.f90 +++ b/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 a/gcc/testsuite/gfortran.dg/goacc/sie.f95 b/gcc/testsuite/gfortran.dg/goacc/sie.f95 index 2d66026..abfe28b 100644 --- a/gcc/testsuite/gfortran.dg/goacc/sie.f95 +++ b/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 a/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 b/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 index b87d26f..5dea42b 100644 --- a/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 +++ b/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 |