From 4c05daa5dda4b17a55372865e74208bb7585820b Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Wed, 23 Jan 2019 02:40:08 -0800 Subject: [PATCH] Make new OpenACC kernels conversion the default; adjust and add tests gcc/c-family/ * c.opt (fopenacc-kernels): Default to "split". gcc/fortran/ * lang.opt (fopenacc-kernels): Default to "split". gcc/ * doc/invoke.texi (-fopenacc-kernels): Update. gcc/testsuite/ * c-c++-common/goacc/note-parallelism-1-kernels-conditional-loop-independent_seq.c: New file. * c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.c: Likewise. * c-c++-common/goacc/note-parallelism-1-kernels-loop-independent_seq.c: Likewise. * c-c++-common/goacc/note-parallelism-1-kernels-loops.c: Likewise. * c-c++-common/goacc/note-parallelism-1-kernels-straight-line.c: Likewise. * c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.c: Likewise. * c-c++-common/goacc/note-parallelism-combined-kernels-loop-independent_seq.c: Likewise. * c-c++-common/goacc/note-parallelism-kernels-conditional-loop-independent_seq.c: Likewise. * c-c++-common/goacc/note-parallelism-kernels-loop-auto.c: Likewise. * c-c++-common/goacc/note-parallelism-kernels-loop-independent_seq.c: Likewise. * c-c++-common/goacc/note-parallelism-kernels-loops.c: Likewise. * c-c++-common/goacc/classify-kernels-unparallelized.c: Update. * 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/dtype-1.c: Likewise. * c-c++-common/goacc/if-clause-2.c: Likewise. * c-c++-common/goacc/kernels-conversion.c: Likewise. * c-c++-common/goacc/kernels-decompose-1.c: Likewise. * c-c++-common/goacc/loop-2-kernels.c: Likewise. * c-c++-common/goacc/note-parallelism.c: Likewise. * c-c++-common/goacc/routine-1.c: Likewise. * c-c++-common/goacc/uninit-dim-clause.c: Likewise. * gfortran.dg/goacc/dtype-1.f95: Likewise. * gfortran.dg/goacc/kernels-conversion.f95: Likewise. * gfortran.dg/goacc/kernels-decompose-1.f95: Likewise. * gfortran.dg/goacc/kernels-tree.f95: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: Update. * testsuite/libgomp.oacc-c-c++-common/avoid-offloading-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/avoid-offloading-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/avoid-offloading-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: Likewise. * testsuite/libgomp.oacc-fortran/avoid-offloading-1.f: Likewise. * testsuite/libgomp.oacc-fortran/avoid-offloading-2.f: Likewise. * testsuite/libgomp.oacc-fortran/avoid-offloading-3.f: Likewise. * testsuite/libgomp.oacc-fortran/initialize_kernels_loops.f90: Likewise. (cherry picked from openacc-gcc-9-branch commit 7035758011d24fbd187a5fdffb59cf42303bdf87) --- gcc/ChangeLog.omp | 4 + gcc/c-family/ChangeLog.omp | 4 + gcc/c-family/c.opt | 2 +- gcc/doc/invoke.texi | 2 +- gcc/fortran/ChangeLog.omp | 4 + gcc/fortran/lang.opt | 2 +- gcc/testsuite/ChangeLog.omp | 39 ++++ .../goacc/classify-kernels-unparallelized.c | 7 +- .../c-c++-common/goacc/classify-kernels.c | 2 +- .../c-c++-common/goacc/classify-parallel.c | 2 +- .../c-c++-common/goacc/classify-routine.c | 2 +- .../c-c++-common/goacc/if-clause-2.c | 1 - .../c-c++-common/goacc/kernels-conversion.c | 10 +- .../c-c++-common/goacc/kernels-decompose-1.c | 1 - .../c-c++-common/goacc/loop-2-kernels.c | 14 +- ...kernels-conditional-loop-independent_seq.c | 129 +++++++++++ .../note-parallelism-1-kernels-loop-auto.c | 126 +++++++++++ ...rallelism-1-kernels-loop-independent_seq.c | 126 +++++++++++ .../goacc/note-parallelism-1-kernels-loops.c | 47 ++++ ...note-parallelism-1-kernels-straight-line.c | 82 +++++++ ...e-parallelism-combined-kernels-loop-auto.c | 121 +++++++++++ ...sm-combined-kernels-loop-independent_seq.c | 121 +++++++++++ ...kernels-conditional-loop-independent_seq.c | 204 ++++++++++++++++++ .../note-parallelism-kernels-loop-auto.c | 138 ++++++++++++ ...parallelism-kernels-loop-independent_seq.c | 138 ++++++++++++ .../goacc/note-parallelism-kernels-loops.c | 50 +++++ .../c-c++-common/goacc/note-parallelism.c | 3 +- gcc/testsuite/c-c++-common/goacc/routine-1.c | 2 +- .../c-c++-common/goacc/uninit-dim-clause.c | 6 +- .../gfortran.dg/goacc/kernels-conversion.f95 | 7 +- .../gfortran.dg/goacc/kernels-decompose-1.f95 | 1 - .../gfortran.dg/goacc/kernels-tree.f95 | 1 - libgomp/ChangeLog.omp | 18 ++ .../kernels-decompose-1.c | 3 +- 34 files changed, 1383 insertions(+), 36 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-conditional-loop-independent_seq.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-independent_seq.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loops.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-straight-line.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-independent_seq.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-conditional-loop-independent_seq.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-auto.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-independent_seq.c create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loops.c diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index 73415e1c5ca3..dda0e6b3409d 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,7 @@ +2019-01-23 Thomas Schwinge + + * doc/invoke.texi (-fopenacc-kernels): Update. + 2019-01-24 Thomas Schwinge * omp-oacc-kernels.c (adjust_region_code_walk_stmt_fn) diff --git a/gcc/c-family/ChangeLog.omp b/gcc/c-family/ChangeLog.omp index e1c4fefcb99e..468c4437b168 100644 --- a/gcc/c-family/ChangeLog.omp +++ b/gcc/c-family/ChangeLog.omp @@ -1,3 +1,7 @@ +2019-01-23 Thomas Schwinge + + * c.opt (fopenacc-kernels): Default to "split". + 2019-01-30 Thomas Schwinge * c.opt (fopenacc-kernels): New flag. diff --git a/gcc/c-family/c.opt b/gcc/c-family/c.opt index 0c0c161fcf49..3f0c0aabf7e7 100644 --- a/gcc/c-family/c.opt +++ b/gcc/c-family/c.opt @@ -1669,7 +1669,7 @@ C ObjC C++ ObjC++ LTO Joined Var(flag_openacc_dims) Specify default OpenACC compute dimensions. fopenacc-kernels= -C ObjC C++ ObjC++ RejectNegative Joined Enum(openacc_kernels) Var(flag_openacc_kernels) Init(OPENACC_KERNELS_PARLOOPS) +C ObjC C++ ObjC++ RejectNegative Joined Enum(openacc_kernels) Var(flag_openacc_kernels) Init(OPENACC_KERNELS_SPLIT) -fopenacc-kernels=[split|parloops] Configure OpenACC 'kernels' constructs handling. Enum diff --git a/gcc/doc/invoke.texi b/gcc/doc/invoke.texi index 1f6d2516c745..77d1ce5e3b69 100644 --- a/gcc/doc/invoke.texi +++ b/gcc/doc/invoke.texi @@ -2195,9 +2195,9 @@ Configure OpenACC 'kernels' constructs handling. With @option{-fopenacc-kernels=split}, OpenACC 'kernels' constructs are split into a sequence of compute constructs, each then handled individually. +This is the default. With @option{-fopenacc-kernels=parloops}, the whole OpenACC 'kernels' constructs is handled by the @samp{parloops} pass. -This is the default. @item -fopenmp @opindex fopenmp diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp index f4d3bd25cf2e..514009825d40 100644 --- a/gcc/fortran/ChangeLog.omp +++ b/gcc/fortran/ChangeLog.omp @@ -1,3 +1,7 @@ +2019-01-23 Thomas Schwinge + + * lang.opt (fopenacc-kernels): Default to "split". + 2019-01-30 Thomas Schwinge * lang.opt (fopenacc-kernels): New flag. diff --git a/gcc/fortran/lang.opt b/gcc/fortran/lang.opt index e6f607da602a..9698a49cd51a 100644 --- a/gcc/fortran/lang.opt +++ b/gcc/fortran/lang.opt @@ -651,7 +651,7 @@ Fortran LTO Joined Var(flag_openacc_dims) ; Documented in C fopenacc-kernels= -Fortran RejectNegative Joined Enum(openacc_kernels) Var(flag_openacc_kernels) Init(OPENACC_KERNELS_PARLOOPS) +Fortran RejectNegative Joined Enum(openacc_kernels) Var(flag_openacc_kernels) Init(OPENACC_KERNELS_SPLIT) ; Documented in C fopenmp diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index 61bc60de4a73..49f1a2bcd10c 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,42 @@ +2019-01-23 Thomas Schwinge + + * c-c++-common/goacc/note-parallelism-1-kernels-conditional-loop-independent_seq.c: + New file. + * c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.c: + Likewise. + * c-c++-common/goacc/note-parallelism-1-kernels-loop-independent_seq.c: + Likewise. + * c-c++-common/goacc/note-parallelism-1-kernels-loops.c: Likewise. + * c-c++-common/goacc/note-parallelism-1-kernels-straight-line.c: + Likewise. + * c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.c: + Likewise. + * c-c++-common/goacc/note-parallelism-combined-kernels-loop-independent_seq.c: + Likewise. + * c-c++-common/goacc/note-parallelism-kernels-conditional-loop-independent_seq.c: + Likewise. + * c-c++-common/goacc/note-parallelism-kernels-loop-auto.c: + Likewise. + * c-c++-common/goacc/note-parallelism-kernels-loop-independent_seq.c: + Likewise. + * c-c++-common/goacc/note-parallelism-kernels-loops.c: Likewise. + * c-c++-common/goacc/classify-kernels-unparallelized.c: Update. + * 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/dtype-1.c: Likewise. + * c-c++-common/goacc/if-clause-2.c: Likewise. + * c-c++-common/goacc/kernels-conversion.c: Likewise. + * c-c++-common/goacc/kernels-decompose-1.c: Likewise. + * c-c++-common/goacc/loop-2-kernels.c: Likewise. + * c-c++-common/goacc/note-parallelism.c: Likewise. + * c-c++-common/goacc/routine-1.c: Likewise. + * c-c++-common/goacc/uninit-dim-clause.c: Likewise. + * gfortran.dg/goacc/dtype-1.f95: Likewise. + * gfortran.dg/goacc/kernels-conversion.f95: Likewise. + * gfortran.dg/goacc/kernels-decompose-1.f95: Likewise. + * gfortran.dg/goacc/kernels-tree.f95: Likewise. + 2019-01-24 Thomas Schwinge * c-c++-common/goacc/kernels-conversion.c: Adjust test. 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 d4c4b2ca237a..9dad2de504cf 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c @@ -1,5 +1,5 @@ /* Check offloaded function's attributes and classification for unparallelized - OpenACC kernels. */ + OpenACC 'kernels'. */ /* { dg-additional-options "-O2" } { dg-additional-options "-fopt-info-optimized-omp" } @@ -13,14 +13,15 @@ extern unsigned int *__restrict a; extern unsigned int *__restrict b; extern unsigned int *__restrict c; -/* An "extern"al mapping of loop iterations/array indices makes the loop - unparallelizable. */ extern unsigned int f (unsigned int); +#pragma acc routine (f) seq void KERNELS () { #pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ for (unsigned int i = 0; i < N; i++) + /* An "extern"al mapping of loop iterations/array indices makes the loop + unparallelizable. */ c[i] = a[f (i)] + b[f (i)]; } diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c index 16e9b9e31d16..f1d461306857 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c @@ -1,5 +1,5 @@ /* Check offloaded function's attributes and classification for OpenACC - kernels. */ + 'kernels'. */ /* { dg-additional-options "-O2" } { dg-additional-options "-fopt-info-optimized-omp" } diff --git a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c index 66a6d1336638..9c80efd65e18 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c @@ -1,5 +1,5 @@ /* Check offloaded function's attributes and classification for OpenACC - parallel. */ + 'parallel'. */ /* { dg-additional-options "-O2" } { dg-additional-options "-fopt-info-optimized-omp" } diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine.c b/gcc/testsuite/c-c++-common/goacc/classify-routine.c index 0b9ba6ea69fc..a4994b061f02 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-routine.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-routine.c @@ -1,5 +1,5 @@ /* Check offloaded function's attributes and classification for OpenACC - routine. */ + 'routine'. */ /* { dg-additional-options "-O2" } { dg-additional-options "-fopt-info-optimized-omp" } diff --git a/gcc/testsuite/c-c++-common/goacc/if-clause-2.c b/gcc/testsuite/c-c++-common/goacc/if-clause-2.c index e17b5dd1107a..9920b4fd175a 100644 --- a/gcc/testsuite/c-c++-common/goacc/if-clause-2.c +++ b/gcc/testsuite/c-c++-common/goacc/if-clause-2.c @@ -1,4 +1,3 @@ -/* { dg-additional-options "-fopenacc-kernels=split" } */ /* { dg-additional-options "-fdump-tree-convert_oacc_kernels" } */ void diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-conversion.c b/gcc/testsuite/c-c++-common/goacc/kernels-conversion.c index ea7eec997fb8..8cb63f00444b 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-conversion.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-conversion.c @@ -1,4 +1,3 @@ -/* { dg-additional-options "-fopenacc-kernels=split" } */ /* { dg-additional-options "-fdump-tree-convert_oacc_kernels" } */ #define N 1024 @@ -52,12 +51,11 @@ main (void) parallelized loop region; and three "old-style" kernel regions. */ /* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single" 1 "convert_oacc_kernels" } } */ /* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_parallelized" 1 "convert_oacc_kernels" } } */ -/* { dg-final { scan-tree-dump-times "oacc_kernels" 3 "convert_oacc_kernels" } } */ +/* { dg-final { scan-tree-dump-times "oacc_kernels " 3 "convert_oacc_kernels" } } */ /* Each of the parallel regions is async, and there is a final call to __builtin_GOACC_wait. */ -/* { dg-final { scan-tree-dump-times "oacc_parallel_kernels.* async\(-1\)" 5 "convert_oacc_kernels" } } */ +/* { dg-final { scan-tree-dump-times "oacc_kernels async\\(-1\\)" 3 "convert_oacc_kernels" } } */ +/* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single async\\(-1\\)" 1 "convert_oacc_kernels" } } */ +/* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_parallelized async\\(-1\\)" 1 "convert_oacc_kernels" } } */ /* { dg-final { scan-tree-dump-times "__builtin_GOACC_wait" 1 "convert_oacc_kernels" } } */ - -/* Check that the original kernels region is removed. */ -/* { dg-final { scan-tree-dump-not "oacc_kernels" "convert_oacc_kernels" } } */ 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 b5d58c372000..7255d692e342 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c @@ -1,6 +1,5 @@ /* Test OpenACC 'kernels' construct decomposition. */ -/* { dg-additional-options "-fopenacc-kernels=split" } */ /* { dg-additional-options "-fopt-info-optimized-omp" } */ /* { dg-additional-options "-O2" } for "parloops". */ diff --git a/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c b/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c index 01515089a7d8..c989222669c0 100644 --- a/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c +++ b/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c @@ -37,7 +37,7 @@ void K(void) for (j = 0; j < 10; j++) { } } -#pragma acc loop seq gang // { dg-error "'seq' overrides" } +#pragma acc loop seq gang // { dg-error "'seq' overrides" "TODO" { xfail *-*-* } } for (i = 0; i < 10; i++) { } @@ -63,7 +63,7 @@ void K(void) for (j = 0; j < 10; j++) { } } -#pragma acc loop seq worker // { dg-error "'seq' overrides" } +#pragma acc loop seq worker // { dg-error "'seq' overrides" "TODO" { xfail *-*-* } } for (i = 0; i < 10; i++) { } #pragma acc loop gang worker @@ -92,7 +92,7 @@ void K(void) for (j = 1; j < 10; j++) { } } -#pragma acc loop seq vector // { dg-error "'seq' overrides" } +#pragma acc loop seq vector // { dg-error "'seq' overrides" "TODO" { xfail *-*-* } } for (i = 0; i < 10; i++) { } #pragma acc loop gang vector @@ -105,7 +105,7 @@ void K(void) #pragma acc loop auto for (i = 0; i < 10; i++) { } -#pragma acc loop seq auto // { dg-error "'seq' overrides" } +#pragma acc loop seq auto // { dg-error "'seq' overrides" "TODO" { xfail *-*-* } } for (i = 0; i < 10; i++) { } #pragma acc loop gang auto // { dg-error "'auto' conflicts" } @@ -147,7 +147,7 @@ void K(void) #pragma acc kernels loop worker(num:5) for (i = 0; i < 10; i++) { } -#pragma acc kernels loop seq worker // { dg-error "'seq' overrides" } +#pragma acc kernels loop seq worker // { dg-error "'seq' overrides" "TODO" { xfail *-*-* } } for (i = 0; i < 10; i++) { } #pragma acc kernels loop gang worker @@ -163,7 +163,7 @@ void K(void) #pragma acc kernels loop vector(length:5) for (i = 0; i < 10; i++) { } -#pragma acc kernels loop seq vector // { dg-error "'seq' overrides" } +#pragma acc kernels loop seq vector // { dg-error "'seq' overrides" "TODO" { xfail *-*-* } } for (i = 0; i < 10; i++) { } #pragma acc kernels loop gang vector @@ -176,7 +176,7 @@ void K(void) #pragma acc kernels loop auto for (i = 0; i < 10; i++) { } -#pragma acc kernels loop seq auto // { dg-error "'seq' overrides" } +#pragma acc kernels loop seq auto // { dg-error "'seq' overrides" "TODO" { xfail *-*-* } } for (i = 0; i < 10; i++) { } #pragma acc kernels loop gang auto // { dg-error "'auto' conflicts" } diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-conditional-loop-independent_seq.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-conditional-loop-independent_seq.c new file mode 100644 index 000000000000..a81d3559dafc --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-conditional-loop-independent_seq.c @@ -0,0 +1,129 @@ +/* Test the output of "-fopt-info-optimized-omp" for an OpenACC 'kernels' + construct containing conditionally executed 'loop' constructs with + 'independent' or 'seq' clauses. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +extern int c; + +int +main () +{ + int x, y, z; + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop seq + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent worker + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent vector + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang vector + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang worker + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent worker vector + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang worker vector + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent worker + for (y = 0; y < 10; y++) +#pragma acc loop independent vector + for (z = 0; z < 10; z++) + ; + +#pragma acc loop independent + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) + ; + +#pragma acc loop independent + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) +#pragma acc loop independent + for (z = 0; z < 10; z++) + ; + +#pragma acc loop seq + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) +#pragma acc loop independent + for (z = 0; z < 10; z++) + ; + +#pragma acc loop independent + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop seq + for (y = 0; y < 10; y++) +#pragma acc loop independent + for (z = 0; z < 10; z++) + ; + +#pragma acc loop independent + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) +#pragma acc loop seq + for (z = 0; z < 10; z++) + ; + +#pragma acc loop seq + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) +#pragma acc loop seq + for (z = 0; z < 10; z++) + ; + } + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.c new file mode 100644 index 000000000000..22ac5399a9d2 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.c @@ -0,0 +1,126 @@ +/* Test the output of "-fopt-info-optimized-omp" for an OpenACC 'kernels' + construct containing 'loop' constructs with explicit or implicit 'auto' + clause. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +int +main () +{ + int x, y, z; + +#pragma acc kernels + /* Strangely indented to keep this similar to other test cases. */ + { +#pragma acc loop /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop auto gang /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop auto worker /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop auto vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop auto gang vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop auto gang worker /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop auto worker vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop auto gang worker vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop auto gang /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto worker + for (y = 0; y < 10; y++) +#pragma acc loop auto vector + for (z = 0; z < 10; z++) + ; + +#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) + ; + +#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop auto + for (z = 0; z < 10; z++) + ; + +#pragma acc loop /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop auto + for (z = 0; z < 10; z++) + ; + +#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop + for (y = 0; y < 10; y++) +#pragma acc loop auto + for (z = 0; z < 10; z++) + ; + +#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop + for (z = 0; z < 10; z++) + ; + +#pragma acc loop /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop + for (z = 0; z < 10; z++) + ; + } + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-independent_seq.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-independent_seq.c new file mode 100644 index 000000000000..a436cd3f007d --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-independent_seq.c @@ -0,0 +1,126 @@ +/* Test the output of "-fopt-info-optimized-omp" for an OpenACC 'kernels' + construct containing 'loop' constructs with 'independent' or 'seq' + clauses. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +int +main () +{ + int x, y, z; + +#pragma acc kernels + /* Strangely indented to keep this similar to other test cases. */ + { +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang vector /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang worker /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent worker vector /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang worker vector /* { dg-message "note: assigned OpenACC gang worker vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) + ; + +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + } + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loops.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loops.c new file mode 100644 index 000000000000..e8b994b5be05 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loops.c @@ -0,0 +1,47 @@ +/* Test the output of "-fopt-info-optimized-omp" for an OpenACC 'kernels' + construct containing loops. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +int +main () +{ + int x, y, z; + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + { + for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + ; + + for (x = 0; x < 10; x++) + ; + + for (x = 0; x < 10; x++) + for (y = 0; y < 10; y++) + for (z = 0; z < 10; z++) + ; + + for (x = 0; x < 10; x++) + ; + + for (x = 0; x < 10; x++) + for (y = 0; y < 10; y++) + ; + + for (x = 0; x < 10; x++) + for (y = 0; y < 10; y++) + for (z = 0; z < 10; z++) + ; + + for (x = 0; x < 10; x++) + for (y = 0; y < 10; y++) + for (z = 0; z < 10; z++) + ; + } + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-straight-line.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-straight-line.c new file mode 100644 index 000000000000..8e40f6217be0 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-straight-line.c @@ -0,0 +1,82 @@ +/* Test the output of "-fopt-info-optimized-omp" for an OpenACC 'kernels' + construct containing straight-line code. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +#pragma acc routine gang +extern int +f_g (int); + +#pragma acc routine worker +extern int +f_w (int); + +#pragma acc routine vector +extern int +f_v (int); + +#pragma acc routine seq +extern int +f_s (int); + +int +main () +{ + int x, y, z; + +#pragma acc kernels /* { dg-warning "region contains gang partitoned code but is not gang partitioned" } */ + { + x = 0; /* { dg-message "note: beginning .gang-single. region in OpenACC .kernels. construct" } */ + y = x < 10; + z = x++; + ; + + y = 0; + z = y < 10; + x -= f_g (y++); /* { dg-message "note: assigned OpenACC gang worker vector loop parallelism" } */ + ; + + x = f_w (0); /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */ + z = f_v (x < 10); /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + y -= f_s (x++); /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + ; + + x = 0; + y = x < 10; + z = (x++); + y = 0; + x = y < 10; + z += (y++); + ; + + x = 0; + y += f_s (x < 10); /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + x++; + y = 0; + y += f_v (y < 10); /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + y++; + z = 0; + y += f_w (z < 10); /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */ + z++; + ; + + x = 0; + y *= f_g ( /* { dg-message "note: assigned OpenACC gang worker vector loop parallelism" } */ + f_w (x < 10) /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */ + + f_g (x < 10) /* { dg-message "note: assigned OpenACC gang worker vector loop parallelism" } */ + ); + x++; + y = 0; + y *= y < 10; + y++; + z = 0; + y *= z < 10; + z++; + ; + } + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.c new file mode 100644 index 000000000000..0254036d7af7 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.c @@ -0,0 +1,121 @@ +/* Test the output of "-fopt-info-optimized-omp" for combined OpenACC 'kernels + loop' constructs with explicit or implicit 'auto' clause. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +int +main () +{ + int x, y, z; + +#pragma acc kernels loop /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop auto gang /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop auto worker /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop auto vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop auto gang vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop auto gang worker /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop auto worker vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop auto gang worker vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop auto gang /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto worker + for (y = 0; y < 10; y++) +#pragma acc loop auto vector + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) + ; + +#pragma acc kernels loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop auto + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop auto + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop + for (y = 0; y < 10; y++) +#pragma acc loop auto + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop + for (z = 0; z < 10; z++) + ; + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-independent_seq.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-independent_seq.c new file mode 100644 index 000000000000..83602a9414d6 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-independent_seq.c @@ -0,0 +1,121 @@ +/* Test the output of "-fopt-info-optimized-omp" for combined OpenACC 'kernels + loop' constructs with 'independent' or 'seq' clauses. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +int +main () +{ + int x, y, z; + +#pragma acc kernels loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent gang vector /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent gang worker /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent worker vector /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent gang worker vector /* { dg-message "note: assigned OpenACC gang worker vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop independent /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) + ; + +#pragma acc kernels loop independent /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-conditional-loop-independent_seq.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-conditional-loop-independent_seq.c new file mode 100644 index 000000000000..e12e0fdae526 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-conditional-loop-independent_seq.c @@ -0,0 +1,204 @@ +/* Test the output of "-fopt-info-optimized-omp" for OpenACC 'kernels' + constructs containing conditionally executed 'loop' constructs with + 'independent' or 'seq' clauses. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +extern int c; + +int +main () +{ + int x, y, z; + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop seq + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent gang + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent worker + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent vector + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent gang vector + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent gang worker + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent worker vector + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent gang worker vector + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent gang + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent worker + for (y = 0; y < 10; y++) +#pragma acc loop independent vector + for (z = 0; z < 10; z++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) +#pragma acc loop independent + for (z = 0; z < 10; z++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop seq + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) +#pragma acc loop independent + for (z = 0; z < 10; z++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop seq + for (y = 0; y < 10; y++) +#pragma acc loop independent + for (z = 0; z < 10; z++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop independent + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) +#pragma acc loop seq + for (z = 0; z < 10; z++) + ; + } + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* Strangely indented to keep this similar to other test cases. */ + if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + { +#pragma acc loop seq + /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent + for (y = 0; y < 10; y++) +#pragma acc loop seq + for (z = 0; z < 10; z++) + ; + } + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-auto.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-auto.c new file mode 100644 index 000000000000..d52b2e860c26 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-auto.c @@ -0,0 +1,138 @@ +/* Test the output of "-fopt-info-optimized-omp" for OpenACC 'kernels' + constructs containing 'loop' constructs with explicit or implicit 'auto' + clause. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +int +main () +{ + int x, y, z; + +#pragma acc kernels +#pragma acc loop /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop auto gang /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop auto worker /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop auto vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop auto gang vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop auto gang worker /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop auto worker vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop auto gang worker vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop auto gang /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto worker + for (y = 0; y < 10; y++) +#pragma acc loop auto vector + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) + ; + +#pragma acc kernels +#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop auto + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop auto + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop + for (y = 0; y < 10; y++) +#pragma acc loop auto + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop auto + for (y = 0; y < 10; y++) +#pragma acc loop + for (z = 0; z < 10; z++) + ; + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-independent_seq.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-independent_seq.c new file mode 100644 index 000000000000..661f7122a2c0 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-independent_seq.c @@ -0,0 +1,138 @@ +/* Test the output of "-fopt-info-optimized-omp" for OpenACC 'kernels' + constructs containing 'loop' constructs with 'independent' or 'seq' + clauses. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +int +main () +{ + int x, y, z; + +#pragma acc kernels +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop independent gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop independent worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop independent vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop independent gang vector /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop independent gang worker /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop independent worker vector /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop independent gang worker vector /* { dg-message "note: assigned OpenACC gang worker vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop independent gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc kernels +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) + ; + +#pragma acc kernels +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */ + for (x = 0; x < 10; x++) +#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loops.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loops.c new file mode 100644 index 000000000000..7587d9d2962b --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loops.c @@ -0,0 +1,50 @@ +/* Test the output of "-fopt-info-optimized-omp" for an OpenACC 'kernels' + construct containing loops. */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +//TODO update accordingly +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +int +main () +{ + int x, y, z; + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + ; + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + ; + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + for (y = 0; y < 10; y++) + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + ; + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + for (y = 0; y < 10; y++) + ; + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + for (y = 0; y < 10; y++) + for (z = 0; z < 10; z++) + ; + +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */ + for (y = 0; y < 10; y++) + for (z = 0; z < 10; z++) + ; + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism.c index 735df7dfad7a..2b49a8be5d12 100644 --- a/gcc/testsuite/c-c++-common/goacc/note-parallelism.c +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism.c @@ -1,4 +1,5 @@ -/* Test the output of "-fopt-info-optimized-omp". */ +/* Test the output of "-fopt-info-optimized-omp" for OpenACC 'parallel' + constructs. */ /* { dg-additional-options "-fopt-info-optimized-omp" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/routine-1.c b/gcc/testsuite/c-c++-common/goacc/routine-1.c index b76c07423059..d68daaaa1dd9 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-1.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-1.c @@ -27,7 +27,7 @@ void nohost (void) int main () { -#pragma acc kernels num_gangs (32) num_workers (32) vector_length (32) +#pragma acc kernels num_gangs (32) num_workers (32) vector_length (32) /* { dg-warning "region contains gang partitoned code but is not gang partitioned" } */ { gang (); worker (); 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 72aacd70f79f..6bc35d777651 100644 --- a/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c +++ b/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c @@ -21,12 +21,12 @@ 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_gangs(i) /* { dg-warning "is used uninitialized in this function" "TODO" { xfail *-*-* } } */ ; - #pragma acc kernels num_workers(j) /* { dg-warning "is used uninitialized in this function" } */ + #pragma acc kernels num_workers(j) /* { dg-warning "is used uninitialized in this function" "TODO" { xfail *-*-* } } */ ; - #pragma acc kernels vector_length(k) /* { dg-warning "is used uninitialized in this function" } */ + #pragma acc kernels vector_length(k) /* { dg-warning "is used uninitialized in this function" "TODO" { xfail *-*-* } } */ ; } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-conversion.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-conversion.f95 index 6604727cf13a..4672d15572ec 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-conversion.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-conversion.f95 @@ -1,4 +1,3 @@ -! { dg-additional-options "-fopenacc-kernels=split" } ! { dg-additional-options "-fdump-tree-convert_oacc_kernels" } program main @@ -50,9 +49,11 @@ end program main ! parallelized loop region; and three "old-style" kernel regions. ! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single" 1 "convert_oacc_kernels" } } ! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_parallelized" 1 "convert_oacc_kernels" } } -! { dg-final { scan-tree-dump-times "oacc_kernels" 3 "convert_oacc_kernels" } } +! { dg-final { scan-tree-dump-times "oacc_kernels " 3 "convert_oacc_kernels" } } ! Each of the parallel regions is async, and there is a final call to ! __builtin_GOACC_wait. -! { dg-final { scan-tree-dump-times "oacc_parallel_kernels.* async\(-1\)" 5 "convert_oacc_kernels" } } +! { dg-final { scan-tree-dump-times "oacc_kernels async\\(-1\\)" 3 "convert_oacc_kernels" } } +! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single async\\(-1\\)" 1 "convert_oacc_kernels" } } +! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_parallelized async\\(-1\\)" 1 "convert_oacc_kernels" } } ! { dg-final { scan-tree-dump-times "__builtin_GOACC_wait" 1 "convert_oacc_kernels" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95 index 520bf034ac6b..8173c3651e16 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95 @@ -1,6 +1,5 @@ ! Test OpenACC 'kernels' construct decomposition. -! { dg-additional-options "-fopenacc-kernels=split" } ! { dg-additional-options "-fopt-info-optimized-omp" } ! { dg-additional-options "-O2" } for "parloops". diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 index b83ca2d8f064..bc9bebac9699 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 @@ -1,6 +1,5 @@ ! { dg-do compile } ! { dg-additional-options "-fdump-tree-original" } -! { dg-additional-options "-fopenacc-kernels=split" } ! { dg-additional-options "-fdump-tree-convert_oacc_kernels" } program test diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index 89360d6f4458..6635ee1110ed 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,21 @@ +2019-01-23 Thomas Schwinge + + * testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: + Update. + * testsuite/libgomp.oacc-c-c++-common/avoid-offloading-1.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/avoid-offloading-2.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/avoid-offloading-3.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: + Likewise. + * testsuite/libgomp.oacc-fortran/avoid-offloading-1.f: Likewise. + * testsuite/libgomp.oacc-fortran/avoid-offloading-2.f: Likewise. + * testsuite/libgomp.oacc-fortran/avoid-offloading-3.f: Likewise. + * testsuite/libgomp.oacc-fortran/initialize_kernels_loops.f90: + Likewise. + 2019-01-24 Thomas Schwinge * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: New diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c index 601e543495e7..bf58a139d3ad 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c @@ -1,4 +1,3 @@ -/* { dg-additional-options "-fopenacc-kernels=split" } */ /* { dg-additional-options "-fopt-info-optimized-omp" } */ #undef NDEBUG @@ -15,7 +14,7 @@ int main() int c = 234; /* { dg-warning "note: beginning .gang-single. region in OpenACC .kernels. construct" } */ #pragma acc loop independent gang /* { dg-warning "note: assigned OpenACC gang loop parallelism" } */ - /* { dg-warning "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } 17 } */ + /* { dg-warning "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } 16 } */ for (int i = 0; i < N; ++i) b[i] = c; -- 2.47.2