diff options
author | Thomas Schwinge <thomas@codesourcery.com> | 2021-04-23 12:23:51 +0200 |
---|---|---|
committer | Thomas Schwinge <thomas@codesourcery.com> | 2021-04-26 12:32:00 +0200 |
commit | 22cff118f7526bec195ed6e41452980820fdf3a8 (patch) | |
tree | b2050907ef4e5e174e850714c98bcdc9f0713562 /gcc | |
parent | 7c640779bf042ceb05fe50260307af88fed75407 (diff) | |
download | gcc-22cff118f7526bec195ed6e41452980820fdf3a8.zip gcc-22cff118f7526bec195ed6e41452980820fdf3a8.tar.gz gcc-22cff118f7526bec195ed6e41452980820fdf3a8.tar.bz2 |
Add '-Wopenacc-parallelism'
... to diagnose potentially suboptimal choices regarding OpenACC parallelism.
Not enabled by default: too noisy ("*potentially* suboptimal choices"); see
XFAILed 'dg-bogus'es.
gcc/c-family/
* c.opt (Wopenacc-parallelism): New.
gcc/fortran/
* lang.opt (Wopenacc-parallelism): New.
gcc/
* omp-offload.c (oacc_validate_dims): Implement
'-Wopenacc-parallelism'.
* doc/invoke.texi (-Wopenacc-parallelism): Document.
gcc/testsuite/
* c-c++-common/goacc/diag-parallelism-1.c: New.
* c-c++-common/goacc/acc-icf.c: Specify '-Wopenacc-parallelism',
and match diagnostics, as appropriate.
* c-c++-common/goacc/classify-kernels-unparallelized.c: Likewise.
* c-c++-common/goacc/classify-kernels.c: Likewise.
* c-c++-common/goacc/classify-parallel.c: Likewise.
* c-c++-common/goacc/classify-routine.c: Likewise.
* c-c++-common/goacc/classify-serial.c: Likewise.
* c-c++-common/goacc/kernels-decompose-1.c: Likewise.
* c-c++-common/goacc/kernels-decompose-2.c: Likewise.
* c-c++-common/goacc/parallel-dims-1.c: Likewise.
* c-c++-common/goacc/parallel-reduction.c: Likewise.
* c-c++-common/goacc/pr70688.c: Likewise.
* c-c++-common/goacc/routine-1.c: Likewise.
* c-c++-common/goacc/routine-level-of-parallelism-2.c: Likewise.
* c-c++-common/goacc/uninit-dim-clause.c: Likewise.
* gfortran.dg/goacc/classify-kernels-unparallelized.f95: Likewise.
* gfortran.dg/goacc/classify-kernels.f95: Likewise.
* gfortran.dg/goacc/classify-parallel.f95: Likewise.
* gfortran.dg/goacc/classify-routine.f95: Likewise.
* gfortran.dg/goacc/classify-serial.f95: Likewise.
* gfortran.dg/goacc/kernels-decompose-1.f95: Likewise.
* gfortran.dg/goacc/kernels-decompose-2.f95: Likewise.
* gfortran.dg/goacc/parallel-tree.f95: Likewise.
* gfortran.dg/goacc/routine-4.f90: Likewise.
* gfortran.dg/goacc/routine-level-of-parallelism-1.f90: Likewise.
* gfortran.dg/goacc/routine-module-mod-1.f90: Likewise.
* gfortran.dg/goacc/routine-multiple-directives-1.f90: Likewise.
* gfortran.dg/goacc/uninit-dim-clause.f95: Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: Specify
'-Wopenacc-parallelism', and match diagnostics, as appropriate.
* testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-w-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/mode-transitions.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/parallel-reduction.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85381-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/private-variables.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/reduction-5.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/reduction-7.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/routine-g-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/routine-w-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/static-variable-1.c:
Likewise.
* testsuite/libgomp.oacc-fortran/optional-private.f90: Likewise.
* testsuite/libgomp.oacc-fortran/par-reduction-2-1.f: Likewise.
* testsuite/libgomp.oacc-fortran/par-reduction-2-2.f: Likewise.
* testsuite/libgomp.oacc-fortran/parallel-dims.f90: Likewise.
* testsuite/libgomp.oacc-fortran/parallel-reduction.f90: Likewise.
* testsuite/libgomp.oacc-fortran/pr84028.f90: Likewise.
* testsuite/libgomp.oacc-fortran/private-variables.f90: Likewise.
* testsuite/libgomp.oacc-fortran/reduction-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/reduction-5.f90: Likewise.
* testsuite/libgomp.oacc-fortran/reduction-6.f90: Likewise.
* testsuite/libgomp.oacc-fortran/routine-7.f90: Likewise.
Co-Authored-By: Nathan Sidwell <nathan@codesourcery.com>
Co-Authored-By: Tom de Vries <vries@codesourcery.com>
Co-Authored-By: Julian Brown <julian@codesourcery.com>
Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com>
Diffstat (limited to 'gcc')
32 files changed, 302 insertions, 1 deletions
diff --git a/gcc/c-family/c.opt b/gcc/c-family/c.opt index 3f8b72c..f1b4c3f 100644 --- a/gcc/c-family/c.opt +++ b/gcc/c-family/c.opt @@ -1037,6 +1037,10 @@ Wold-style-definition C ObjC Var(warn_old_style_definition) Init(-1) Warning Warn if an old-style parameter definition is used. +Wopenacc-parallelism +C C++ Var(warn_openacc_parallelism) Warning +Warn about potentially suboptimal choices related to OpenACC parallelism. + Wopenmp-simd C C++ Var(warn_openmp_simd) Warning LangEnabledBy(C C++,Wall) Warn if a simd directive is overridden by the vectorizer cost model. diff --git a/gcc/doc/invoke.texi b/gcc/doc/invoke.texi index e98b096..40cacc6 100644 --- a/gcc/doc/invoke.texi +++ b/gcc/doc/invoke.texi @@ -364,7 +364,9 @@ Objective-C and Objective-C++ Dialects}. -Wmissing-include-dirs -Wmissing-noreturn -Wno-missing-profile @gol -Wno-multichar -Wmultistatement-macros -Wnonnull -Wnonnull-compare @gol -Wnormalized=@r{[}none@r{|}id@r{|}nfc@r{|}nfkc@r{]} @gol --Wnull-dereference -Wno-odr -Wopenmp-simd @gol +-Wnull-dereference -Wno-odr @gol +-Wopenacc-parallelism @gol +-Wopenmp-simd @gol -Wno-overflow -Woverlength-strings -Wno-override-init-side-effects @gol -Wpacked -Wno-packed-bitfield-compat -Wpacked-not-aligned -Wpadded @gol -Wparentheses -Wno-pedantic-ms-format @gol @@ -8749,6 +8751,12 @@ Do not warn about compile-time overflow in constant expressions. Warn about One Definition Rule violations during link-time optimization. Enabled by default. +@item -Wopenacc-parallelism +@opindex Wopenacc-parallelism +@opindex Wno-openacc-parallelism +@cindex OpenACC accelerator programming +Warn about potentially suboptimal choices related to OpenACC parallelism. + @item -Wopenmp-simd @opindex Wopenmp-simd @opindex Wno-openmp-simd diff --git a/gcc/fortran/lang.opt b/gcc/fortran/lang.opt index 388ef8c..6db01c7 100644 --- a/gcc/fortran/lang.opt +++ b/gcc/fortran/lang.opt @@ -285,6 +285,10 @@ Wuse-without-only Fortran Var(warn_use_without_only) Warning Warn about USE statements that have no ONLY qualifier. +Wopenacc-parallelism +Fortran +; Documented in C + Wopenmp-simd Fortran ; Documented in C diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c index ba0937f..1612461 100644 --- a/gcc/omp-offload.c +++ b/gcc/omp-offload.c @@ -928,6 +928,35 @@ oacc_validate_dims (tree fn, tree attrs, int *dims, int level, unsigned used) pos = TREE_CHAIN (pos); } + bool check = true; +#ifdef ACCEL_COMPILER + check = false; +#endif + if (check + && warn_openacc_parallelism + && !lookup_attribute ("oacc kernels", DECL_ATTRIBUTES (fn))) + { + static char const *const axes[] = + /* Must be kept in sync with GOMP_DIM enumeration. */ + { "gang", "worker", "vector" }; + for (ix = level >= 0 ? level : 0; ix != GOMP_DIM_MAX; ix++) + if (dims[ix] < 0) + ; /* Defaulting axis. */ + else if ((used & GOMP_DIM_MASK (ix)) && dims[ix] == 1) + /* There is partitioned execution, but the user requested a + dimension size of 1. They're probably confused. */ + warning_at (DECL_SOURCE_LOCATION (fn), OPT_Wopenacc_parallelism, + "region contains %s partitioned code but" + " is not %s partitioned", axes[ix], axes[ix]); + else if (!(used & GOMP_DIM_MASK (ix)) && dims[ix] != 1) + /* The dimension is explicitly partitioned to non-unity, but + no use is made within the region. */ + warning_at (DECL_SOURCE_LOCATION (fn), OPT_Wopenacc_parallelism, + "region is %s partitioned but" + " does not contain %s partitioned code", + axes[ix], axes[ix]); + } + bool changed = targetm.goacc.validate_dims (fn, dims, level, used); /* Default anything left to 1 or a partitioned default. */ diff --git a/gcc/testsuite/c-c++-common/goacc/acc-icf.c b/gcc/testsuite/c-c++-common/goacc/acc-icf.c index 98b536c..9cf119b 100644 --- a/gcc/testsuite/c-c++-common/goacc/acc-icf.c +++ b/gcc/testsuite/c-c++-common/goacc/acc-icf.c @@ -2,7 +2,12 @@ /* { dg-additional-options "-fopenacc -O2 -fdump-ipa-icf" } */ +/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting + aspects of that functionality. */ + #pragma acc routine gang +/* { dg-bogus "warning: region is worker partitioned but does not contain worker partitioned code" "TODO default 'gang' 'vector'" { xfail *-*-* } .+3 } + TODO It's the compiler's own decision to not use 'worker' parallelism here, so it doesn't make sense to bother the user about it. */ int routine1 (int n) { @@ -16,6 +21,8 @@ routine1 (int n) } #pragma acc routine gang +/* { dg-bogus "warning: region is worker partitioned but does not contain worker partitioned code" "TODO default 'gang' 'vector'" { xfail *-*-* } .+3 } + TODO It's the compiler's own decision to not use 'worker' parallelism here, so it doesn't make sense to bother the user about it. */ int routine2 (int n) { diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c index d4c4b2c..218f624 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c @@ -7,6 +7,9 @@ { dg-additional-options "-fdump-tree-parloops1-all" } { dg-additional-options "-fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting + aspects of that functionality. */ + #define N 1024 extern unsigned int *__restrict a; diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c index 16e9b9e..95a150c 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c @@ -7,6 +7,9 @@ { dg-additional-options "-fdump-tree-parloops1-all" } { dg-additional-options "-fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting + aspects of that functionality. */ + #define N 1024 extern unsigned int *__restrict a; diff --git a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c index 933d766..230e70c 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c @@ -6,6 +6,9 @@ { dg-additional-options "-fdump-tree-ompexp" } { dg-additional-options "-fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting + aspects of that functionality. */ + #define N 1024 extern unsigned int *__restrict a; diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine.c b/gcc/testsuite/c-c++-common/goacc/classify-routine.c index 0b9ba6e..81fe369 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-routine.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-routine.c @@ -6,6 +6,9 @@ { dg-additional-options "-fdump-tree-ompexp" } { dg-additional-options "-fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting + aspects of that functionality. */ + #define N 1024 extern unsigned int *__restrict a; diff --git a/gcc/testsuite/c-c++-common/goacc/classify-serial.c b/gcc/testsuite/c-c++-common/goacc/classify-serial.c index 94ace1b..ae052ae 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-serial.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-serial.c @@ -6,6 +6,9 @@ { dg-additional-options "-fdump-tree-ompexp" } { dg-additional-options "-fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting + aspects of that functionality. */ + #define N 1024 extern unsigned int *__restrict a; @@ -15,6 +18,11 @@ extern unsigned int *__restrict c; void SERIAL () { #pragma acc serial loop copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */ + /* { dg-bogus "warning: region contains gang partitioned code but is not gang partitioned" "TODO 'serial'" { xfail *-*-* } .-1 } + { dg-bogus "warning: region contains worker partitioned code but is not worker partitioned" "" { target *-*-* } .-2 } + { dg-bogus "warning: region contains vector partitioned code but is not vector partitioned" "TODO 'serial'" { xfail *-*-* } .-3 } + TODO Should we really diagnose this if the user explicitly requested 'serial'? + TODO Should we instead diagnose ('-Wextra' category?) that the user may enable use of parallelism if replacing 'serial' with 'parallel', if applicable? */ for (unsigned int i = 0; i < N; i++) c[i] = a[i] + b[i]; } diff --git a/gcc/testsuite/c-c++-common/goacc/diag-parallelism-1.c b/gcc/testsuite/c-c++-common/goacc/diag-parallelism-1.c new file mode 100644 index 0000000..6ee7bd0 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/diag-parallelism-1.c @@ -0,0 +1,124 @@ +/* Diagnostics about potentially suboptimal choices related to OpenACC + parallelism. + + { dg-additional-options "-Wopenacc-parallelism" } +*/ + + +//TODO 'kernels' + +//TODO 'serial' + +//TODO 'routine' + +//TODO Fortran + + +static void f1 () +{ + int ary[10]; + + +#pragma acc parallel num_gangs (1) + /* { dg-warning "region contains gang partitioned code but is not gang partitioned" "" { target *-*-* } .-1 } */ + { + #pragma acc loop gang + for (int i = 0; i < 10; i++) + ary[i] = i; + } + +#pragma acc parallel num_workers (1) + /* { dg-warning "region contains worker partitioned code but is not worker partitioned" "" { target *-*-* } .-1 } */ + { + #pragma acc loop worker + for (int i = 0; i < 10; i++) + ary[i] = i; + } + +#pragma acc parallel vector_length (1) + /* { dg-warning "region contains vector partitioned code but is not vector partitioned" "" { target *-*-* } .-1 } */ + { + #pragma acc loop vector + for (int i = 0; i < 10; i++) + ary[i] = i; + } +} + + +static void f2 () +{ + int ary[10]; + + +#pragma acc parallel num_gangs (8) + /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .-1 } */ + { + #pragma acc loop worker + for (int i = 0; i < 10; i++) + ary[i] = i; + } + +#pragma acc parallel num_gangs (8) + /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .-1 } */ + { + #pragma acc loop vector + for (int i = 0; i < 10; i++) + ary[i] = i; + } + +#pragma acc parallel num_gangs (8) + /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .-1 } */ + { + #pragma acc loop worker vector + for (int i = 0; i < 10; i++) + ary[i] = i; + } + +#pragma acc parallel num_workers (8) + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } */ + { + #pragma acc loop gang + for (int i = 0; i < 10; i++) + ary[i] = i; + } + +#pragma acc parallel num_workers (8) + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } */ + { + #pragma acc loop vector + for (int i = 0; i < 10; i++) + ary[i] = i; + } + +#pragma acc parallel num_workers (8) + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } */ + { + #pragma acc loop gang vector + for (int i = 0; i < 10; i++) + ary[i] = i; + } + +#pragma acc parallel vector_length (8) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } */ + { + #pragma acc loop gang + for (int i = 0; i < 10; i++) + ary[i] = i; + } + +#pragma acc parallel vector_length (8) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } */ + { + #pragma acc loop worker + for (int i = 0; i < 10; i++) + ary[i] = i; + } + +#pragma acc parallel vector_length (8) + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } */ + { + #pragma acc loop gang worker + for (int i = 0; i < 10; i++) + ary[i] = i; + } +} diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c index 87219c8..f549cba 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c @@ -5,6 +5,9 @@ /* { dg-additional-options "--param=openacc-kernels=decompose" } { dg-additional-options "-fdump-tree-omp_oacc_kernels_decompose" } */ +/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting + aspects of that functionality. */ + /* See also '../../gfortran.dg/goacc/kernels-decompose-1.f95'. */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c index 3781e75..cdf85d4 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c @@ -4,6 +4,9 @@ /* { dg-additional-options "--param=openacc-kernels=decompose" } /* { dg-additional-options "-O2" } for 'parloops'. */ +/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting + aspects of that functionality. */ + /* See also '../../gfortran.dg/goacc/kernels-decompose-2.f95'. */ /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' @@ -116,6 +119,7 @@ main () } #pragma acc kernels + /* { dg-bogus "warning: region contains gang partitioned code but is not gang partitioned" "TODO 'kernels'" { xfail *-*-* } .-1 } */ { y = f_g (a[5]); /* { dg-line l_part[incr c_part] } */ /*TODO If such a construct is placed in its own part (like it is, here), can't this actually use gang paralelism, instead of "gang-single"? 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 57f682f..2a8d35d 100644 --- a/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c +++ b/gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c @@ -1,11 +1,18 @@ /* Valid use of OpenACC parallelism dimensions clauses: num_gangs, num_workers, vector_length. */ +/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting + aspects of that functionality. */ + 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) + /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO runtime" { xfail *-*-* } .-1 } + { dg-bogus "warning: region is worker partitioned but does not contain worker partitioned code" "TODO runtime" { xfail *-*-* } .-2 } + { dg-bogus "warning: region is vector partitioned but does not contain vector partitioned code" "TODO runtime" { xfail *-*-* } .-3 } + TODO 'region is [...] partitioned' isn't correct for 'i == 1'. */ ; } diff --git a/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c b/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c index d7cc947..c5c0edc 100644 --- a/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c +++ b/gcc/testsuite/c-c++-common/goacc/parallel-reduction.c @@ -1,3 +1,6 @@ +/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting + aspects of that functionality. */ + int main () { @@ -7,6 +10,7 @@ main () #pragma acc data copy (dummy) { #pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum) + /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 } */ { int v = 5; sum += 10 + v; diff --git a/gcc/testsuite/c-c++-common/goacc/pr70688.c b/gcc/testsuite/c-c++-common/goacc/pr70688.c index 5a23665..e96a853 100644 --- a/gcc/testsuite/c-c++-common/goacc/pr70688.c +++ b/gcc/testsuite/c-c++-common/goacc/pr70688.c @@ -1,3 +1,6 @@ +/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting + aspects of that functionality. */ + const int n = 100; int @@ -22,6 +25,7 @@ parallel_reduction () #pragma acc data copy (dummy) { #pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum) + /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 } */ { int v = 5; sum += 10 + v; @@ -37,10 +41,12 @@ main () int i, s = 0; #pragma acc parallel num_gangs (10) copy (s) reduction (+:s) + /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 } */ for (i = 0; i < n; i++) s += i+1; #pragma acc parallel num_gangs (10) reduction (+:s) copy (s) + /* { dg-bogus "warning: region is gang partitioned but does not contain gang partitioned code" "TODO 'reduction'" { xfail *-*-* } .-1 } */ for (i = 0; i < n; i++) s += i+1; diff --git a/gcc/testsuite/c-c++-common/goacc/routine-1.c b/gcc/testsuite/c-c++-common/goacc/routine-1.c index a756922..051f793 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-1.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-1.c @@ -1,15 +1,23 @@ +/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting + aspects of that functionality. */ #pragma acc routine gang +/* { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .+3 } + { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .+2 } + { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .+1 } */ void gang (void) { } #pragma acc routine worker +/* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .+2 } + { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .+1 } */ void worker (void) { } #pragma acc routine vector +/* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .+1 } */ void vector (void) { } diff --git a/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c index a066f2b..33678fe 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c @@ -2,7 +2,13 @@ with the OpenACC 'routine' directive. The Fortran counterpart is '../../gfortran.dg/goacc/routine-level-of-parallelism-1.f90'. */ +/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting + aspects of that functionality. */ + #pragma acc routine gang +/* { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .+3 } + { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .+2 } + { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .+1 } */ void g_1 (void) { } 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 827dac7..d9b143b 100644 --- a/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c +++ b/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c @@ -1,16 +1,22 @@ /* { dg-additional-options "-Wuninitialized" } */ +/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting + aspects of that functionality. */ + void acc_parallel() { int i, j, k; #pragma acc parallel num_gangs(i) /* { dg-warning "is used uninitialized" } */ + /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .-1 } */ ; #pragma acc parallel num_workers(j) /* { dg-warning "is used uninitialized" } */ + /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } */ ; #pragma acc parallel vector_length(k) /* { dg-warning "is used uninitialized" } */ + /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } */ ; } diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 index 6cca3d6..cb5251a 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 @@ -7,6 +7,9 @@ ! { dg-additional-options "-fdump-tree-parloops1-all" } ! { dg-additional-options "-fdump-tree-oaccdevlow" } +! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting +! aspects of that functionality. + program main implicit none integer, parameter :: n = 1024 diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 index 715a983..07aaf06 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 @@ -7,6 +7,9 @@ ! { dg-additional-options "-fdump-tree-parloops1-all" } ! { dg-additional-options "-fdump-tree-oaccdevlow" } +! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting +! aspects of that functionality. + program main implicit none integer, parameter :: n = 1024 diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 index 01f06bb..a41e0e6 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 @@ -6,6 +6,9 @@ ! { dg-additional-options "-fdump-tree-ompexp" } ! { dg-additional-options "-fdump-tree-oaccdevlow" } +! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting +! aspects of that functionality. + program main implicit none integer, parameter :: n = 1024 diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 index 401d527..52cc870 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 @@ -6,6 +6,9 @@ ! { dg-additional-options "-fdump-tree-ompexp" } ! { dg-additional-options "-fdump-tree-oaccdevlow" } +! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting +! aspects of that functionality. + subroutine ROUTINE !$acc routine worker integer, parameter :: n = 1024 diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 index 51061af..6d402e6 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 @@ -6,6 +6,9 @@ ! { dg-additional-options "-fdump-tree-ompexp" } ! { dg-additional-options "-fdump-tree-oaccdevlow" } +! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting +! aspects of that functionality. + program main implicit none integer, parameter :: n = 1024 @@ -15,6 +18,9 @@ program main call setup(a, b) !$acc serial loop copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } + ! { dg-bogus "warning: region contains gang partitioned code but is not gang partitioned" "TODO 'serial'" { xfail *-*-* } .-1 } + ! { dg-bogus "warning: region contains worker partitioned code but is not worker partitioned" "" { target *-*-* } .-2 } + ! { dg-bogus "warning: region contains vector partitioned code but is not vector partitioned" "TODO 'serial'" { xfail *-*-* } .-3 } do i = 0, n - 1 c(i) = a(i) + b(i) end do diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95 index e252350..ddaf7f8 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95 @@ -5,6 +5,9 @@ ! { dg-additional-options "--param=openacc-kernels=decompose" } ! { dg-additional-options "-fdump-tree-omp_oacc_kernels_decompose" } +! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting +! aspects of that functionality. + ! See also '../../c-c++-common/goacc/kernels-decompose-1.c'. ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95 index cc12b77..a3cec83 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95 @@ -4,6 +4,9 @@ ! { dg-additional-options "--param=openacc-kernels=decompose" } ! { dg-additional-options "-O2" } for 'parloops'. +! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting +! aspects of that functionality. + ! See also '../../c-c++-common/goacc/kernels-decompose-2.c'. ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' @@ -119,6 +122,7 @@ program main !$acc end kernels !$acc kernels + ! { dg-bogus "warning: region contains gang partitioned code but is not gang partitioned" "TODO 'kernels'" { xfail *-*-* } .-1 } y = f_g (a(5)) ! { dg-line l_part[incr c_part] } !TODO If such a construct is placed in its own part (like it is, here), can't this actually use gang paralelism, instead of "gang-single"? ! { dg-message "note: beginning 'gang-single' part in OpenACC 'kernels' region" "" { target *-*-* } l_part$c_part } diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 index e33653b..c51cb2a 100644 --- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 @@ -2,6 +2,9 @@ ! test for tree-dump-original and spaces-commas +! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting +! aspects of that functionality. + program test implicit none integer :: q, i, j, k, m, n, o, p, r, s, t, u, v, w @@ -12,6 +15,9 @@ program test !$acc no_create(n) & !$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) & !$acc deviceptr(u), private(v), firstprivate(w) + ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .-1 } + ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 } + ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 } !$acc end parallel end program test diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-4.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-4.f90 index 6714c7b..53b1fbe 100644 --- a/gcc/testsuite/gfortran.dg/goacc/routine-4.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/routine-4.f90 @@ -1,5 +1,8 @@ ! Test invalid calls to routines. +! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting +! aspects of that functionality. + module param integer, parameter :: N = 32 end module param @@ -120,6 +123,9 @@ contains subroutine gang (a) ! { dg-message "declared here" 3 } !$acc routine gang + ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .-2 } + ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-3 } + ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-4 } integer, intent (inout) :: a(N) integer :: i @@ -130,6 +136,8 @@ contains subroutine worker (a) ! { dg-message "declared here" 2 } !$acc routine worker + ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 } + ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 } integer, intent (inout) :: a(N) integer :: i @@ -140,6 +148,7 @@ contains subroutine vector (a) ! { dg-message "declared here" } !$acc routine vector + ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } integer, intent (inout) :: a(N) integer :: i diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90 index 83b8c24..99929f1 100644 --- a/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90 @@ -2,8 +2,14 @@ ! with the OpenACC routine directive. The C/C++ counterpart is ! '../../c-c++-common/goacc/routine-level-of-parallelism-2.c'. +! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting +! aspects of that functionality. + subroutine g_1 !$acc routine gang + ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .-2 } + ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-3 } + ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-4 } end subroutine g_1 subroutine s_1_2a diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90 index 23c673f..2210650 100644 --- a/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/routine-module-mod-1.f90 @@ -2,6 +2,9 @@ ! { dg-additional-options "-fopt-info-optimized-omp" } +! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting +! aspects of that functionality. + module routine_module_mod_1 contains subroutine s_1 @@ -53,6 +56,7 @@ contains subroutine g_1 implicit none !$acc routine gang + ! { dg-bogus "warning: region is worker partitioned but does not contain worker partitioned code" "TODO default 'gang' 'vector'" { xfail *-*-* } .-3 } integer :: i diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90 index 6e12ee9..622a9d9 100644 --- a/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90 @@ -1,5 +1,8 @@ ! Check for valid cases of multiple OpenACC 'routine' directives. +! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting +! aspects of that functionality. + SUBROUTINE s_1 !$ACC ROUTINE(s_1) !$ACC ROUTINE(s_1) SEQ @@ -17,12 +20,14 @@ !$ACC ROUTINE VECTOR !$ACC ROUTINE(v_1) VECTOR !$ACC ROUTINE VECTOR +! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-5 } END SUBROUTINE v_1 SUBROUTINE v_2 !$ACC ROUTINE(v_2) VECTOR !$ACC ROUTINE VECTOR !$ACC ROUTINE(v_2) VECTOR +! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-4 } END SUBROUTINE v_2 SUBROUTINE sub_1 diff --git a/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 b/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 index c77d47a..59e923f 100644 --- a/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 @@ -1,16 +1,22 @@ ! { dg-additional-options "-Wuninitialized" } +! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting +! aspects of that functionality. + subroutine acc_parallel implicit none integer :: i, j, k !$acc parallel num_gangs(i) ! { dg-warning "is used uninitialized" } + ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .-1 } !$acc end parallel !$acc parallel num_workers(j) ! { dg-warning "is used uninitialized" } + ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } !$acc end parallel !$acc parallel vector_length(k) ! { dg-warning "is used uninitialized" } + ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-1 } !$acc end parallel end subroutine acc_parallel |