* [PATCH 0/3] Add OpenACC diagnostics to -fopt-info-note-omp
@ 2018-07-25 15:29 Cesar Philippidis
2018-07-25 15:29 ` [PATCH 1/3] Correct the reported line number in fortran combined OpenACC directives Cesar Philippidis
` (2 more replies)
0 siblings, 3 replies; 12+ messages in thread
From: Cesar Philippidis @ 2018-07-25 15:29 UTC (permalink / raw)
To: gcc-patches, jakub
This patch series extends -fopt-info-note-omp to include OpenACC loop
diagnostics when it is used in conjunction with -fopenacc. At present,
the diagnostics are limited to reporting how OpenACC loops are
partitioned, e.g., seq, gang, worker or vector. The major advantage of
this diagnostics is that it informs the user how GCC automatically
partitions independent loops, i.e., acc loops without any parallelism
clauses inside acc parallel regions. This information provides the
user with insights on how to select num_gangs, num_workers and
vector_length for their application.
All three patches in this series are independent from one
another. Patches 1 and 2 fix diagnostics bugs involving incorrect line
numbers. Patch 3 is responsible for generating the actual diagnostics.
Cesar
^ permalink raw reply [flat|nested] 12+ messages in thread* [PATCH 1/3] Correct the reported line number in fortran combined OpenACC directives 2018-07-25 15:29 [PATCH 0/3] Add OpenACC diagnostics to -fopt-info-note-omp Cesar Philippidis @ 2018-07-25 15:29 ` Cesar Philippidis 2018-07-25 15:32 ` Marek Polacek 2018-07-25 15:29 ` [PATCH 3/3] Add user-friendly OpenACC diagnostics regarding detected parallelism Cesar Philippidis 2018-07-25 15:29 ` [PATCH 2/3] Correct the reported line number in c++ combined OpenACC directives Cesar Philippidis 2 siblings, 1 reply; 12+ messages in thread From: Cesar Philippidis @ 2018-07-25 15:29 UTC (permalink / raw) To: gcc-patches, jakub The fortran FE incorrectly records the line locations of combined acc loop directives when it lowers the construct to gimple. Usually this isn't a problem because the fortran FE is able to report problems with acc loops itself. However, there will be inaccuracies if the ME tries to use those locations. Note that test cases are inconspicuously absent in this patch. However, without this bug fix, -fopt-info-note-omp will report bogus line numbers. This code patch will be tested in a later patch in this series. Is this OK for trunk? I bootstrapped and regtested it on x86_64 with nvptx offloading. Thanks, Cesar 2018-XX-YY Cesar Philippidis <cesar@codesourcery.com> gcc/fortran/ * trans-openmp.c (gfc_trans_oacc_combined_directive): Set the location of combined acc loops. (cherry picked from gomp-4_0-branch r245653) diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index f038f4c..e7707d0 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -3869,6 +3869,7 @@ gfc_trans_oacc_combined_directive (gfc_code *code) gfc_omp_clauses construct_clauses, loop_clauses; tree stmt, oacc_clauses = NULL_TREE; enum tree_code construct_code; + location_t loc = input_location; switch (code->op) { @@ -3930,12 +3931,16 @@ gfc_trans_oacc_combined_directive (gfc_code *code) else pushlevel (); stmt = gfc_trans_omp_do (code, EXEC_OACC_LOOP, pblock, &loop_clauses, NULL); + + if (CAN_HAVE_LOCATION_P (stmt)) + SET_EXPR_LOCATION (stmt, loc); + if (TREE_CODE (stmt) != BIND_EXPR) stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0)); else poplevel (0, 0); - stmt = build2_loc (input_location, construct_code, void_type_node, stmt, - oacc_clauses); + + stmt = build2_loc (loc, construct_code, void_type_node, stmt, oacc_clauses); gfc_add_expr_to_block (&block, stmt); return gfc_finish_block (&block); } -- 2.7.4 ^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [PATCH 1/3] Correct the reported line number in fortran combined OpenACC directives 2018-07-25 15:29 ` [PATCH 1/3] Correct the reported line number in fortran combined OpenACC directives Cesar Philippidis @ 2018-07-25 15:32 ` Marek Polacek 2018-07-25 15:53 ` Cesar Philippidis 0 siblings, 1 reply; 12+ messages in thread From: Marek Polacek @ 2018-07-25 15:32 UTC (permalink / raw) To: Cesar Philippidis; +Cc: gcc-patches, jakub On Wed, Jul 25, 2018 at 08:29:17AM -0700, Cesar Philippidis wrote: > The fortran FE incorrectly records the line locations of combined acc > loop directives when it lowers the construct to gimple. Usually this > isn't a problem because the fortran FE is able to report problems with > acc loops itself. However, there will be inaccuracies if the ME tries > to use those locations. > > Note that test cases are inconspicuously absent in this patch. > However, without this bug fix, -fopt-info-note-omp will report bogus > line numbers. This code patch will be tested in a later patch in > this series. > > Is this OK for trunk? I bootstrapped and regtested it on x86_64 with > nvptx offloading. > > Thanks, > Cesar > > 2018-XX-YY Cesar Philippidis <cesar@codesourcery.com> > > gcc/fortran/ > * trans-openmp.c (gfc_trans_oacc_combined_directive): Set the > location of combined acc loops. > > (cherry picked from gomp-4_0-branch r245653) > > diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c > index f038f4c..e7707d0 100644 > --- a/gcc/fortran/trans-openmp.c > +++ b/gcc/fortran/trans-openmp.c > @@ -3869,6 +3869,7 @@ gfc_trans_oacc_combined_directive (gfc_code *code) > gfc_omp_clauses construct_clauses, loop_clauses; > tree stmt, oacc_clauses = NULL_TREE; > enum tree_code construct_code; > + location_t loc = input_location; > > switch (code->op) > { > @@ -3930,12 +3931,16 @@ gfc_trans_oacc_combined_directive (gfc_code *code) > else > pushlevel (); > stmt = gfc_trans_omp_do (code, EXEC_OACC_LOOP, pblock, &loop_clauses, NULL); > + > + if (CAN_HAVE_LOCATION_P (stmt)) > + SET_EXPR_LOCATION (stmt, loc); This is protected_set_expr_location. Marek ^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [PATCH 1/3] Correct the reported line number in fortran combined OpenACC directives 2018-07-25 15:32 ` Marek Polacek @ 2018-07-25 15:53 ` Cesar Philippidis 2018-12-09 13:07 ` Thomas Schwinge 0 siblings, 1 reply; 12+ messages in thread From: Cesar Philippidis @ 2018-07-25 15:53 UTC (permalink / raw) To: Marek Polacek; +Cc: gcc-patches, jakub [-- Attachment #1: Type: text/plain, Size: 1930 bytes --] On 07/25/2018 08:32 AM, Marek Polacek wrote: > On Wed, Jul 25, 2018 at 08:29:17AM -0700, Cesar Philippidis wrote: >> The fortran FE incorrectly records the line locations of combined acc >> loop directives when it lowers the construct to gimple. Usually this >> isn't a problem because the fortran FE is able to report problems with >> acc loops itself. However, there will be inaccuracies if the ME tries >> to use those locations. >> >> Note that test cases are inconspicuously absent in this patch. >> However, without this bug fix, -fopt-info-note-omp will report bogus >> line numbers. This code patch will be tested in a later patch in >> this series. >> >> Is this OK for trunk? I bootstrapped and regtested it on x86_64 with >> nvptx offloading. >> >> Thanks, >> Cesar >> >> 2018-XX-YY Cesar Philippidis <cesar@codesourcery.com> >> >> gcc/fortran/ >> * trans-openmp.c (gfc_trans_oacc_combined_directive): Set the >> location of combined acc loops. >> >> (cherry picked from gomp-4_0-branch r245653) >> >> diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c >> index f038f4c..e7707d0 100644 >> --- a/gcc/fortran/trans-openmp.c >> +++ b/gcc/fortran/trans-openmp.c >> @@ -3869,6 +3869,7 @@ gfc_trans_oacc_combined_directive (gfc_code *code) >> gfc_omp_clauses construct_clauses, loop_clauses; >> tree stmt, oacc_clauses = NULL_TREE; >> enum tree_code construct_code; >> + location_t loc = input_location; >> >> switch (code->op) >> { >> @@ -3930,12 +3931,16 @@ gfc_trans_oacc_combined_directive (gfc_code *code) >> else >> pushlevel (); >> stmt = gfc_trans_omp_do (code, EXEC_OACC_LOOP, pblock, &loop_clauses, NULL); >> + >> + if (CAN_HAVE_LOCATION_P (stmt)) >> + SET_EXPR_LOCATION (stmt, loc); > > This is protected_set_expr_location. Neat, thanks! This patch includes that correction. Is it ok for trunk after bootstrapping and regression testing? Thanks, Cesar [-- Attachment #2: 0001-Correct-the-reported-line-number-in-fortran-combined.patch --] [-- Type: text/x-patch, Size: 1403 bytes --] 2018-XX-YY Cesar Philippidis <cesar@codesourcery.com> gcc/fortran/ * trans-openmp.c (gfc_trans_oacc_combined_directive): Set the location of combined acc loops. (cherry picked from gomp-4_0-branch r245653) --- gcc/fortran/trans-openmp.c | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index f038f4c5bf8..b549c682533 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -3869,6 +3869,7 @@ gfc_trans_oacc_combined_directive (gfc_code *code) gfc_omp_clauses construct_clauses, loop_clauses; tree stmt, oacc_clauses = NULL_TREE; enum tree_code construct_code; + location_t loc = input_location; switch (code->op) { @@ -3929,13 +3930,16 @@ gfc_trans_oacc_combined_directive (gfc_code *code) pblock = █ else pushlevel (); + stmt = gfc_trans_omp_do (code, EXEC_OACC_LOOP, pblock, &loop_clauses, NULL); + protected_set_expr_location (stmt, loc); + if (TREE_CODE (stmt) != BIND_EXPR) stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0)); else poplevel (0, 0); - stmt = build2_loc (input_location, construct_code, void_type_node, stmt, - oacc_clauses); + + stmt = build2_loc (loc, construct_code, void_type_node, stmt, oacc_clauses); gfc_add_expr_to_block (&block, stmt); return gfc_finish_block (&block); } -- 2.17.1 ^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [PATCH 1/3] Correct the reported line number in fortran combined OpenACC directives 2018-07-25 15:53 ` Cesar Philippidis @ 2018-12-09 13:07 ` Thomas Schwinge 0 siblings, 0 replies; 12+ messages in thread From: Thomas Schwinge @ 2018-12-09 13:07 UTC (permalink / raw) To: gcc-patches, fortran; +Cc: Marek Polacek, jakub Hi! On Wed, 25 Jul 2018 08:53:35 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote: > On 07/25/2018 08:32 AM, Marek Polacek wrote: > > On Wed, Jul 25, 2018 at 08:29:17AM -0700, Cesar Philippidis wrote: > >> The fortran FE incorrectly records the line locations of combined acc > >> loop directives when it lowers the construct to gimple. After a bit of preparational work to "use existing middle end checking for Fortran OpenACC loop clauses"... > >> Usually this > >> isn't a problem because the fortran FE is able to report problems with > >> acc loops itself. ..., the Fortran front end is no longer doing that, and... > >> However, there will be inaccuracies if the ME tries > >> to use those locations. > >> > >> Note that test cases are inconspicuously absent in this patch. ..., I've been able to verify your changes by translating your C++ test case into Fortran. > >> However, without this bug fix, -fopt-info-note-omp will report bogus > >> line numbers. This code patch will be tested in a later patch in > >> this series. > >> + if (CAN_HAVE_LOCATION_P (stmt)) > >> + SET_EXPR_LOCATION (stmt, loc); > > > > This is protected_set_expr_location. > > Neat, thanks! This patch includes that correction. Is it ok for trunk > after bootstrapping and regression testing? Thanks, committed to trunk in r266924: commit a43ff24656fa8224b249e159ea81e629ffa32664 Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> Date: Sun Dec 9 12:49:20 2018 +0000 Correct the reported line number in Fortran combined OpenACC directives gcc/fortran/ * trans-openmp.c (gfc_trans_oacc_combined_directive): Set the location of combined acc loops. gcc/testsuite/ * gfortran.dg/goacc/combined-directives-3.f90: New file. Reviewed-by: Thomas Schwinge <thomas@codesourcery.com> git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@266924 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/fortran/ChangeLog | 5 +++++ gcc/fortran/trans-openmp.c | 5 +++-- gcc/testsuite/ChangeLog | 4 ++++ .../c-c++-common/goacc/combined-directives-3.c | 1 + .../gfortran.dg/goacc/combined-directives-3.f90 | 26 ++++++++++++++++++++++ 5 files changed, 39 insertions(+), 2 deletions(-) diff --git gcc/fortran/ChangeLog gcc/fortran/ChangeLog index c6eb05174f69..e74bda7a1362 100644 --- gcc/fortran/ChangeLog +++ gcc/fortran/ChangeLog @@ -1,3 +1,8 @@ +2018-12-09 Cesar Philippidis <cesar@codesourcery.com> + + * trans-openmp.c (gfc_trans_oacc_combined_directive): Set the + location of combined acc loops. + 2018-12-09 Thomas Schwinge <thomas@codesourcery.com> * openmp.c (resolve_oacc_loop_blocks): Remove checking of OpenACC diff --git gcc/fortran/trans-openmp.c gcc/fortran/trans-openmp.c index c9fc4e49c450..bf3f46939e39 100644 --- gcc/fortran/trans-openmp.c +++ gcc/fortran/trans-openmp.c @@ -3878,6 +3878,7 @@ gfc_trans_oacc_combined_directive (gfc_code *code) gfc_omp_clauses construct_clauses, loop_clauses; tree stmt, oacc_clauses = NULL_TREE; enum tree_code construct_code; + location_t loc = input_location; switch (code->op) { @@ -3939,12 +3940,12 @@ gfc_trans_oacc_combined_directive (gfc_code *code) else pushlevel (); stmt = gfc_trans_omp_do (code, EXEC_OACC_LOOP, pblock, &loop_clauses, NULL); + protected_set_expr_location (stmt, loc); if (TREE_CODE (stmt) != BIND_EXPR) stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0)); else poplevel (0, 0); - stmt = build2_loc (input_location, construct_code, void_type_node, stmt, - oacc_clauses); + stmt = build2_loc (loc, construct_code, void_type_node, stmt, oacc_clauses); gfc_add_expr_to_block (&block, stmt); return gfc_finish_block (&block); } diff --git gcc/testsuite/ChangeLog gcc/testsuite/ChangeLog index 6b26f6f510db..19bc532c9d57 100644 --- gcc/testsuite/ChangeLog +++ gcc/testsuite/ChangeLog @@ -1,3 +1,7 @@ +2018-12-09 Thomas Schwinge <thomas@codesourcery.com> + + * gfortran.dg/goacc/combined-directives-3.f90: New file. + 2018-12-09 Cesar Philippidis <cesar@codesourcery.com> * c-c++-common/goacc/combined-directives-3.c: New test. diff --git gcc/testsuite/c-c++-common/goacc/combined-directives-3.c gcc/testsuite/c-c++-common/goacc/combined-directives-3.c index 77d418262eac..c6e31c26a8f1 100644 --- gcc/testsuite/c-c++-common/goacc/combined-directives-3.c +++ gcc/testsuite/c-c++-common/goacc/combined-directives-3.c @@ -1,5 +1,6 @@ /* Verify the accuracy of the line number associated with combined constructs. */ +/* See also "../../gfortran.dg/goacc/combined-directives-3.f90". */ int main () diff --git gcc/testsuite/gfortran.dg/goacc/combined-directives-3.f90 gcc/testsuite/gfortran.dg/goacc/combined-directives-3.f90 new file mode 100644 index 000000000000..b138822827f6 --- /dev/null +++ gcc/testsuite/gfortran.dg/goacc/combined-directives-3.f90 @@ -0,0 +1,26 @@ +! Verify the accuracy of the line number associated with combined constructs. +! See "../../c-c++-common/goacc/combined-directives-3.c". + +subroutine test + implicit none + integer x, y, z + + !$acc parallel loop seq auto ! { dg-error "'seq' overrides other OpenACC loop specifiers" } + do x = 0, 10 + !$acc loop + do y = 0, 10 + end do + end do + !$acc end parallel loop + + !$acc parallel loop gang auto ! { dg-error "'auto' conflicts with other OpenACC loop specifiers" } + do x = 0, 10 + !$acc loop worker auto ! { dg-error "'auto' conflicts with other OpenACC loop specifiers" } + do y = 0, 10 + !$acc loop vector + do z = 0, 10 + end do + end do + end do + !$acc end parallel loop +end subroutine test Grüße Thomas ^ permalink raw reply [flat|nested] 12+ messages in thread
* [PATCH 3/3] Add user-friendly OpenACC diagnostics regarding detected parallelism. 2018-07-25 15:29 [PATCH 0/3] Add OpenACC diagnostics to -fopt-info-note-omp Cesar Philippidis 2018-07-25 15:29 ` [PATCH 1/3] Correct the reported line number in fortran combined OpenACC directives Cesar Philippidis @ 2018-07-25 15:29 ` Cesar Philippidis 2018-07-26 8:34 ` Richard Biener 2018-07-25 15:29 ` [PATCH 2/3] Correct the reported line number in c++ combined OpenACC directives Cesar Philippidis 2 siblings, 1 reply; 12+ messages in thread From: Cesar Philippidis @ 2018-07-25 15:29 UTC (permalink / raw) To: gcc-patches, jakub This patch teaches GCC to inform the user how it assigned parallelism to each OpenACC loop at compile time using the -fopt-info-note-omp flag. For instance, given the acc parallel loop nest: #pragma acc parallel loop for (...) #pragma acc loop vector for (...) GCC will report somthing like foo.c:4:0: note: Detected parallelism <acc loop gang worker> foo.c:6:0: note: Detected parallelism <acc loop vector> Note how only the inner loop specifies vector parallelism. In this example, GCC automatically assigned gang and worker parallelism to the outermost loop. Perhaps, going forward, it would be useful to distinguish which parallelism was specified by the user and which was assigned by the compiler. But that can be added in a follow up patch. Is this patch OK for trunk? I bootstrapped and regtested it for x86_64 with nvptx offloading. Thanks, Cesar 2018-XX-YY Cesar Philippidis <cesar@codesourcery.com> gcc/ * omp-offload.c (inform_oacc_loop): New function. (execute_oacc_device_lower): Use it to display loop parallelism. gcc/testsuite/ * c-c++-common/goacc/note-parallelism.c: New test. * gfortran.dg/goacc/note-parallelism.f90: New test. (cherry picked from gomp-4_0-branch r245683, and gcc/testsuite/ parts of r245770) diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c index 0abf028..66b99bb 100644 --- a/gcc/omp-offload.c +++ b/gcc/omp-offload.c @@ -866,6 +866,31 @@ debug_oacc_loop (oacc_loop *loop) dump_oacc_loop (stderr, loop, 0); } +/* Provide diagnostics on OpenACC loops LOOP, its siblings and its + children. */ + +static void +inform_oacc_loop (oacc_loop *loop) +{ + const char *seq = loop->mask == 0 ? " seq" : ""; + const char *gang = loop->mask & GOMP_DIM_MASK (GOMP_DIM_GANG) + ? " gang" : ""; + const char *worker = loop->mask & GOMP_DIM_MASK (GOMP_DIM_WORKER) + ? " worker" : ""; + const char *vector = loop->mask & GOMP_DIM_MASK (GOMP_DIM_VECTOR) + ? " vector" : ""; + dump_location_t loc = dump_location_t::from_location_t (loop->loc); + + dump_printf_loc (MSG_NOTE, loc, + "Detected parallelism <acc loop%s%s%s%s>\n", seq, gang, + worker, vector); + + if (loop->child) + inform_oacc_loop (loop->child); + if (loop->sibling) + inform_oacc_loop (loop->sibling); +} + /* DFS walk of basic blocks BB onwards, creating OpenACC loop structures as we go. By construction these loops are properly nested. */ @@ -1533,6 +1558,8 @@ execute_oacc_device_lower () dump_oacc_loop (dump_file, loops, 0); fprintf (dump_file, "\n"); } + if (dump_enabled_p () && loops->child) + inform_oacc_loop (loops->child); /* Offloaded targets may introduce new basic blocks, which require dominance information to update SSA. */ diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism.c new file mode 100644 index 0000000..3ec794c --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism.c @@ -0,0 +1,61 @@ +/* Test the output of -fopt-info-note-omp. */ + +/* { dg-additional-options "-fopt-info-note-omp" } */ + +int +main () +{ + int x, y, z; + +#pragma acc parallel loop seq /* { dg-message "note: Detected parallelism <acc loop seq>" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop gang /* { dg-message "note: Detected parallelism <acc loop gang>" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop worker /* { dg-message "note: Detected parallelism <acc loop worker>" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop vector /* { dg-message "note: Detected parallelism <acc loop vector>" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop gang vector /* { dg-message "note: Detected parallelism <acc loop gang vector>" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop gang worker /* { dg-message "note: Detected parallelism <acc loop gang worker>" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop worker vector /* { dg-message "note: Detected parallelism <acc loop worker vector>" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop gang worker vector /* { dg-message "note: Detected parallelism <acc loop gang worker vector>" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop /* { dg-message "note: Detected parallelism <acc loop gang vector>" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop /* { dg-message "note: Detected parallelism <acc loop gang worker>" } */ + for (x = 0; x < 10; x++) +#pragma acc loop /* { dg-message "note: Detected parallelism <acc loop vector>" } */ + for (y = 0; y < 10; y++) + ; + +#pragma acc parallel loop gang /* { dg-message "note: Detected parallelism <acc loop gang>" } */ + for (x = 0; x < 10; x++) +#pragma acc loop worker /* { dg-message "note: Detected parallelism <acc loop worker>" } */ + for (y = 0; y < 10; y++) +#pragma acc loop vector /* { dg-message "note: Detected parallelism <acc loop vector>" } */ + for (z = 0; z < 10; z++) + ; + + return 0; +} diff --git a/gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 b/gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 new file mode 100644 index 0000000..a0c78c5 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 @@ -0,0 +1,62 @@ +! Test the output of -fopt-info-note-omp. + +! { dg-additional-options "-fopt-info-note-omp" } + +program test + implicit none + + integer x, y, z + + !$acc parallel loop seq ! { dg-message "note: Detected parallelism <acc loop seq>" } + do x = 1, 10 + end do + + !$acc parallel loop gang ! { dg-message "note: Detected parallelism <acc loop gang>" } + do x = 1, 10 + end do + + !$acc parallel loop worker ! { dg-message "note: Detected parallelism <acc loop worker>" } + do x = 1, 10 + end do + + !$acc parallel loop vector ! { dg-message "note: Detected parallelism <acc loop vector>" } + do x = 1, 10 + end do + + !$acc parallel loop gang vector ! { dg-message "note: Detected parallelism <acc loop gang vector>" } + do x = 1, 10 + end do + + !$acc parallel loop gang worker ! { dg-message "note: Detected parallelism <acc loop gang worker>" } + do x = 1, 10 + end do + + !$acc parallel loop worker vector ! { dg-message "note: Detected parallelism <acc loop worker vector>" } + do x = 1, 10 + end do + + !$acc parallel loop gang worker vector ! { dg-message "note: Detected parallelism <acc loop gang worker vector>" } + do x = 1, 10 + end do + + !$acc parallel loop ! { dg-message "note: Detected parallelism <acc loop gang vector>" } + do x = 1, 10 + end do + + !$acc parallel loop ! { dg-message "note: Detected parallelism <acc loop gang worker>" } + do x = 1, 10 + !$acc loop ! { dg-message "note: Detected parallelism <acc loop vector>" } + do y = 1, 10 + end do + end do + + !$acc parallel loop gang ! { dg-message "note: Detected parallelism <acc loop gang>" } + do x = 1, 10 + !$acc loop worker ! { dg-message "note: Detected parallelism <acc loop worker>" } + do y = 1, 10 + !$acc loop vector ! { dg-message "note: Detected parallelism <acc loop vector>" } + do z = 1, 10 + end do + end do + end do +end program test -- 2.7.4 ^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [PATCH 3/3] Add user-friendly OpenACC diagnostics regarding detected parallelism. 2018-07-25 15:29 ` [PATCH 3/3] Add user-friendly OpenACC diagnostics regarding detected parallelism Cesar Philippidis @ 2018-07-26 8:34 ` Richard Biener 2018-07-26 14:14 ` Cesar Philippidis 0 siblings, 1 reply; 12+ messages in thread From: Richard Biener @ 2018-07-26 8:34 UTC (permalink / raw) To: cesar; +Cc: GCC Patches, Jakub Jelinek On Wed, Jul 25, 2018 at 5:30 PM Cesar Philippidis <cesar@codesourcery.com> wrote: > > This patch teaches GCC to inform the user how it assigned parallelism > to each OpenACC loop at compile time using the -fopt-info-note-omp > flag. For instance, given the acc parallel loop nest: > > #pragma acc parallel loop > for (...) > #pragma acc loop vector > for (...) > > GCC will report somthing like > > foo.c:4:0: note: Detected parallelism <acc loop gang worker> > foo.c:6:0: note: Detected parallelism <acc loop vector> > > Note how only the inner loop specifies vector parallelism. In this > example, GCC automatically assigned gang and worker parallelism to the > outermost loop. Perhaps, going forward, it would be useful to > distinguish which parallelism was specified by the user and which was > assigned by the compiler. But that can be added in a follow up patch. > > Is this patch OK for trunk? I bootstrapped and regtested it for x86_64 > with nvptx offloading. Shouldn't this use MSG_OPTIMIZED_LOCATIONS instead? Are there any other optinfo notes emitted? Like when despite pragmas loops are not handled or so? > Thanks, > Cesar > > 2018-XX-YY Cesar Philippidis <cesar@codesourcery.com> > > gcc/ > * omp-offload.c (inform_oacc_loop): New function. > (execute_oacc_device_lower): Use it to display loop parallelism. > > gcc/testsuite/ > * c-c++-common/goacc/note-parallelism.c: New test. > * gfortran.dg/goacc/note-parallelism.f90: New test. > > (cherry picked from gomp-4_0-branch r245683, and gcc/testsuite/ parts of > r245770) > > diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c > index 0abf028..66b99bb 100644 > --- a/gcc/omp-offload.c > +++ b/gcc/omp-offload.c > @@ -866,6 +866,31 @@ debug_oacc_loop (oacc_loop *loop) > dump_oacc_loop (stderr, loop, 0); > } > > +/* Provide diagnostics on OpenACC loops LOOP, its siblings and its > + children. */ > + > +static void > +inform_oacc_loop (oacc_loop *loop) > +{ > + const char *seq = loop->mask == 0 ? " seq" : ""; > + const char *gang = loop->mask & GOMP_DIM_MASK (GOMP_DIM_GANG) > + ? " gang" : ""; > + const char *worker = loop->mask & GOMP_DIM_MASK (GOMP_DIM_WORKER) > + ? " worker" : ""; > + const char *vector = loop->mask & GOMP_DIM_MASK (GOMP_DIM_VECTOR) > + ? " vector" : ""; > + dump_location_t loc = dump_location_t::from_location_t (loop->loc); > + > + dump_printf_loc (MSG_NOTE, loc, > + "Detected parallelism <acc loop%s%s%s%s>\n", seq, gang, > + worker, vector); > + > + if (loop->child) > + inform_oacc_loop (loop->child); > + if (loop->sibling) > + inform_oacc_loop (loop->sibling); > +} > + > /* DFS walk of basic blocks BB onwards, creating OpenACC loop > structures as we go. By construction these loops are properly > nested. */ > @@ -1533,6 +1558,8 @@ execute_oacc_device_lower () > dump_oacc_loop (dump_file, loops, 0); > fprintf (dump_file, "\n"); > } > + if (dump_enabled_p () && loops->child) > + inform_oacc_loop (loops->child); > > /* Offloaded targets may introduce new basic blocks, which require > dominance information to update SSA. */ > diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism.c > new file mode 100644 > index 0000000..3ec794c > --- /dev/null > +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism.c > @@ -0,0 +1,61 @@ > +/* Test the output of -fopt-info-note-omp. */ > + > +/* { dg-additional-options "-fopt-info-note-omp" } */ > + > +int > +main () > +{ > + int x, y, z; > + > +#pragma acc parallel loop seq /* { dg-message "note: Detected parallelism <acc loop seq>" } */ > + for (x = 0; x < 10; x++) > + ; > + > +#pragma acc parallel loop gang /* { dg-message "note: Detected parallelism <acc loop gang>" } */ > + for (x = 0; x < 10; x++) > + ; > + > +#pragma acc parallel loop worker /* { dg-message "note: Detected parallelism <acc loop worker>" } */ > + for (x = 0; x < 10; x++) > + ; > + > +#pragma acc parallel loop vector /* { dg-message "note: Detected parallelism <acc loop vector>" } */ > + for (x = 0; x < 10; x++) > + ; > + > +#pragma acc parallel loop gang vector /* { dg-message "note: Detected parallelism <acc loop gang vector>" } */ > + for (x = 0; x < 10; x++) > + ; > + > +#pragma acc parallel loop gang worker /* { dg-message "note: Detected parallelism <acc loop gang worker>" } */ > + for (x = 0; x < 10; x++) > + ; > + > +#pragma acc parallel loop worker vector /* { dg-message "note: Detected parallelism <acc loop worker vector>" } */ > + for (x = 0; x < 10; x++) > + ; > + > +#pragma acc parallel loop gang worker vector /* { dg-message "note: Detected parallelism <acc loop gang worker vector>" } */ > + for (x = 0; x < 10; x++) > + ; > + > +#pragma acc parallel loop /* { dg-message "note: Detected parallelism <acc loop gang vector>" } */ > + for (x = 0; x < 10; x++) > + ; > + > +#pragma acc parallel loop /* { dg-message "note: Detected parallelism <acc loop gang worker>" } */ > + for (x = 0; x < 10; x++) > +#pragma acc loop /* { dg-message "note: Detected parallelism <acc loop vector>" } */ > + for (y = 0; y < 10; y++) > + ; > + > +#pragma acc parallel loop gang /* { dg-message "note: Detected parallelism <acc loop gang>" } */ > + for (x = 0; x < 10; x++) > +#pragma acc loop worker /* { dg-message "note: Detected parallelism <acc loop worker>" } */ > + for (y = 0; y < 10; y++) > +#pragma acc loop vector /* { dg-message "note: Detected parallelism <acc loop vector>" } */ > + for (z = 0; z < 10; z++) > + ; > + > + return 0; > +} > diff --git a/gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 b/gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 > new file mode 100644 > index 0000000..a0c78c5 > --- /dev/null > +++ b/gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 > @@ -0,0 +1,62 @@ > +! Test the output of -fopt-info-note-omp. > + > +! { dg-additional-options "-fopt-info-note-omp" } > + > +program test > + implicit none > + > + integer x, y, z > + > + !$acc parallel loop seq ! { dg-message "note: Detected parallelism <acc loop seq>" } > + do x = 1, 10 > + end do > + > + !$acc parallel loop gang ! { dg-message "note: Detected parallelism <acc loop gang>" } > + do x = 1, 10 > + end do > + > + !$acc parallel loop worker ! { dg-message "note: Detected parallelism <acc loop worker>" } > + do x = 1, 10 > + end do > + > + !$acc parallel loop vector ! { dg-message "note: Detected parallelism <acc loop vector>" } > + do x = 1, 10 > + end do > + > + !$acc parallel loop gang vector ! { dg-message "note: Detected parallelism <acc loop gang vector>" } > + do x = 1, 10 > + end do > + > + !$acc parallel loop gang worker ! { dg-message "note: Detected parallelism <acc loop gang worker>" } > + do x = 1, 10 > + end do > + > + !$acc parallel loop worker vector ! { dg-message "note: Detected parallelism <acc loop worker vector>" } > + do x = 1, 10 > + end do > + > + !$acc parallel loop gang worker vector ! { dg-message "note: Detected parallelism <acc loop gang worker vector>" } > + do x = 1, 10 > + end do > + > + !$acc parallel loop ! { dg-message "note: Detected parallelism <acc loop gang vector>" } > + do x = 1, 10 > + end do > + > + !$acc parallel loop ! { dg-message "note: Detected parallelism <acc loop gang worker>" } > + do x = 1, 10 > + !$acc loop ! { dg-message "note: Detected parallelism <acc loop vector>" } > + do y = 1, 10 > + end do > + end do > + > + !$acc parallel loop gang ! { dg-message "note: Detected parallelism <acc loop gang>" } > + do x = 1, 10 > + !$acc loop worker ! { dg-message "note: Detected parallelism <acc loop worker>" } > + do y = 1, 10 > + !$acc loop vector ! { dg-message "note: Detected parallelism <acc loop vector>" } > + do z = 1, 10 > + end do > + end do > + end do > +end program test > -- > 2.7.4 > ^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [PATCH 3/3] Add user-friendly OpenACC diagnostics regarding detected parallelism. 2018-07-26 8:34 ` Richard Biener @ 2018-07-26 14:14 ` Cesar Philippidis 2018-12-14 21:03 ` Add user-friendly diagnostics for OpenACC loop parallelism assigned (was: [PATCH 3/3] Add user-friendly OpenACC diagnostics regarding detected parallelism) Thomas Schwinge 0 siblings, 1 reply; 12+ messages in thread From: Cesar Philippidis @ 2018-07-26 14:14 UTC (permalink / raw) To: Richard Biener; +Cc: GCC Patches, Jakub Jelinek [-- Attachment #1: Type: text/plain, Size: 2018 bytes --] On 07/26/2018 01:33 AM, Richard Biener wrote: > On Wed, Jul 25, 2018 at 5:30 PM Cesar Philippidis > <cesar@codesourcery.com> wrote: >> >> This patch teaches GCC to inform the user how it assigned parallelism >> to each OpenACC loop at compile time using the -fopt-info-note-omp >> flag. For instance, given the acc parallel loop nest: >> >> #pragma acc parallel loop >> for (...) >> #pragma acc loop vector >> for (...) >> >> GCC will report somthing like >> >> foo.c:4:0: note: Detected parallelism <acc loop gang worker> >> foo.c:6:0: note: Detected parallelism <acc loop vector> >> >> Note how only the inner loop specifies vector parallelism. In this >> example, GCC automatically assigned gang and worker parallelism to the >> outermost loop. Perhaps, going forward, it would be useful to >> distinguish which parallelism was specified by the user and which was >> assigned by the compiler. But that can be added in a follow up patch. >> >> Is this patch OK for trunk? I bootstrapped and regtested it for x86_64 >> with nvptx offloading. > > Shouldn't this use MSG_OPTIMIZED_LOCATIONS instead? Are there > any other optinfo notes emitted? Like when despite pragmas loops > are not handled or so? Early on I was just using the diagnostics in omp-grid.c as a model, but yes, it does make sense to use MSG_OPTIMIZED_LOCATIONS instead of MSG_NOTE. And no, these are the only optinfo notes that we're emitting at the moment. All of the other diagnostics are just errors and warnings, although we probably should revisit that for some of the forthcoming acc routine diagnostics. Going forward, now that there's in interest in automatic parallelism inside acc kernels, we do plan on expanding the diagnostics. The attached revised patch now uses MSG_OPTIMIZED_LOCATIONS for the diagnostics. If this gets approved for trunk, I'll go ahead and backport it to og8 and update the OpenACC wiki to change the usage of -fopt-info-note-omp to -fopt-info-optimized-omp. Is this OK for trunk? Thanks, Cesar [-- Attachment #2: 0003-Add-user-friendly-OpenACC-diagnostics-regarding-dete.patch --] [-- Type: text/x-patch, Size: 6884 bytes --] 2018-XX-YY Cesar Philippidis <cesar@codesourcery.com> gcc/ * omp-offload.c (inform_oacc_loop): New function. (execute_oacc_device_lower): Use it to display loop parallelism. gcc/testsuite/ * c-c++-common/goacc/note-parallelism.c: New test. * gfortran.dg/goacc/note-parallelism.f90: New test. (cherry picked from gomp-4_0-branch r245683, and gcc/testsuite/ parts of r245770) use MSG_OPTIMIZED_LOCATIONS instead of MSG_NOTE --- gcc/omp-offload.c | 27 ++++++++ .../c-c++-common/goacc/note-parallelism.c | 61 ++++++++++++++++++ .../gfortran.dg/goacc/note-parallelism.f90 | 62 +++++++++++++++++++ 3 files changed, 150 insertions(+) create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism.c create mode 100644 gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c index 0abf0283c9e..3582dda3d1a 100644 --- a/gcc/omp-offload.c +++ b/gcc/omp-offload.c @@ -866,6 +866,31 @@ debug_oacc_loop (oacc_loop *loop) dump_oacc_loop (stderr, loop, 0); } +/* Provide diagnostics on OpenACC loops LOOP, its siblings and its + children. */ + +static void +inform_oacc_loop (oacc_loop *loop) +{ + const char *seq = loop->mask == 0 ? " seq" : ""; + const char *gang = loop->mask & GOMP_DIM_MASK (GOMP_DIM_GANG) + ? " gang" : ""; + const char *worker = loop->mask & GOMP_DIM_MASK (GOMP_DIM_WORKER) + ? " worker" : ""; + const char *vector = loop->mask & GOMP_DIM_MASK (GOMP_DIM_VECTOR) + ? " vector" : ""; + dump_location_t loc = dump_location_t::from_location_t (loop->loc); + + dump_printf_loc (MSG_OPTIMIZED_LOCATIONS, loc, + "Detected parallelism <acc loop%s%s%s%s>\n", seq, gang, + worker, vector); + + if (loop->child) + inform_oacc_loop (loop->child); + if (loop->sibling) + inform_oacc_loop (loop->sibling); +} + /* DFS walk of basic blocks BB onwards, creating OpenACC loop structures as we go. By construction these loops are properly nested. */ @@ -1533,6 +1558,8 @@ execute_oacc_device_lower () dump_oacc_loop (dump_file, loops, 0); fprintf (dump_file, "\n"); } + if (dump_enabled_p () && loops->child) + inform_oacc_loop (loops->child); /* Offloaded targets may introduce new basic blocks, which require dominance information to update SSA. */ diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism.c new file mode 100644 index 00000000000..2e50d86cd23 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism.c @@ -0,0 +1,61 @@ +/* Test the output of -fopt-info-note-omp. */ + +/* { dg-additional-options "-fopt-info-note-optimized" } */ + +int +main () +{ + int x, y, z; + +#pragma acc parallel loop seq /* { dg-message "note: Detected parallelism <acc loop seq>" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop gang /* { dg-message "note: Detected parallelism <acc loop gang>" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop worker /* { dg-message "note: Detected parallelism <acc loop worker>" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop vector /* { dg-message "note: Detected parallelism <acc loop vector>" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop gang vector /* { dg-message "note: Detected parallelism <acc loop gang vector>" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop gang worker /* { dg-message "note: Detected parallelism <acc loop gang worker>" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop worker vector /* { dg-message "note: Detected parallelism <acc loop worker vector>" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop gang worker vector /* { dg-message "note: Detected parallelism <acc loop gang worker vector>" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop /* { dg-message "note: Detected parallelism <acc loop gang vector>" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop /* { dg-message "note: Detected parallelism <acc loop gang worker>" } */ + for (x = 0; x < 10; x++) +#pragma acc loop /* { dg-message "note: Detected parallelism <acc loop vector>" } */ + for (y = 0; y < 10; y++) + ; + +#pragma acc parallel loop gang /* { dg-message "note: Detected parallelism <acc loop gang>" } */ + for (x = 0; x < 10; x++) +#pragma acc loop worker /* { dg-message "note: Detected parallelism <acc loop worker>" } */ + for (y = 0; y < 10; y++) +#pragma acc loop vector /* { dg-message "note: Detected parallelism <acc loop vector>" } */ + for (z = 0; z < 10; z++) + ; + + return 0; +} diff --git a/gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 b/gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 new file mode 100644 index 00000000000..2029abfa939 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 @@ -0,0 +1,62 @@ +! Test the output of -fopt-info-note-omp. + +! { dg-additional-options "-fopt-info-note-optimized" } + +program test + implicit none + + integer x, y, z + + !$acc parallel loop seq ! { dg-message "note: Detected parallelism <acc loop seq>" } + do x = 1, 10 + end do + + !$acc parallel loop gang ! { dg-message "note: Detected parallelism <acc loop gang>" } + do x = 1, 10 + end do + + !$acc parallel loop worker ! { dg-message "note: Detected parallelism <acc loop worker>" } + do x = 1, 10 + end do + + !$acc parallel loop vector ! { dg-message "note: Detected parallelism <acc loop vector>" } + do x = 1, 10 + end do + + !$acc parallel loop gang vector ! { dg-message "note: Detected parallelism <acc loop gang vector>" } + do x = 1, 10 + end do + + !$acc parallel loop gang worker ! { dg-message "note: Detected parallelism <acc loop gang worker>" } + do x = 1, 10 + end do + + !$acc parallel loop worker vector ! { dg-message "note: Detected parallelism <acc loop worker vector>" } + do x = 1, 10 + end do + + !$acc parallel loop gang worker vector ! { dg-message "note: Detected parallelism <acc loop gang worker vector>" } + do x = 1, 10 + end do + + !$acc parallel loop ! { dg-message "note: Detected parallelism <acc loop gang vector>" } + do x = 1, 10 + end do + + !$acc parallel loop ! { dg-message "note: Detected parallelism <acc loop gang worker>" } + do x = 1, 10 + !$acc loop ! { dg-message "note: Detected parallelism <acc loop vector>" } + do y = 1, 10 + end do + end do + + !$acc parallel loop gang ! { dg-message "note: Detected parallelism <acc loop gang>" } + do x = 1, 10 + !$acc loop worker ! { dg-message "note: Detected parallelism <acc loop worker>" } + do y = 1, 10 + !$acc loop vector ! { dg-message "note: Detected parallelism <acc loop vector>" } + do z = 1, 10 + end do + end do + end do +end program test -- 2.17.1 ^ permalink raw reply [flat|nested] 12+ messages in thread
* Add user-friendly diagnostics for OpenACC loop parallelism assigned (was: [PATCH 3/3] Add user-friendly OpenACC diagnostics regarding detected parallelism) 2018-07-26 14:14 ` Cesar Philippidis @ 2018-12-14 21:03 ` Thomas Schwinge 2019-01-28 21:51 ` [og8] Add user-friendly diagnostics for OpenACC loop parallelism assigned Thomas Schwinge 0 siblings, 1 reply; 12+ messages in thread From: Thomas Schwinge @ 2018-12-14 21:03 UTC (permalink / raw) To: gcc-patches; +Cc: Jakub Jelinek, Richard Biener Hi! On Thu, 26 Jul 2018 07:14:21 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote: > On 07/26/2018 01:33 AM, Richard Biener wrote: > > On Wed, Jul 25, 2018 at 5:30 PM Cesar Philippidis > > <cesar@codesourcery.com> wrote: > >> > >> This patch Thanks! > >> teaches GCC to inform the user how it assigned parallelism > >> to each OpenACC loop at compile time Hence, I changed that diagnostig to "assigned OpenACC [...] loop parallelism" instead of "Detected parallelism <acc loop [...]>". > >> using the -fopt-info-note-omp > >> flag. For instance, given the acc parallel loop nest: > >> > >> #pragma acc parallel loop > >> for (...) > >> #pragma acc loop vector > >> for (...) > >> > >> GCC will report somthing like > >> > >> foo.c:4:0: note: Detected parallelism <acc loop gang worker> > >> foo.c:6:0: note: Detected parallelism <acc loop vector> > >> > >> Note how only the inner loop specifies vector parallelism. In this > >> example, GCC automatically assigned gang and worker parallelism to the > >> outermost loop. Perhaps, going forward, it would be useful to > >> distinguish which parallelism was specified by the user and which was > >> assigned by the compiler. But that can be added in a follow up patch. ACK. > The attached revised patch now uses MSG_OPTIMIZED_LOCATIONS for the > diagnostics. > Is this OK for trunk? I further changed that to make it build ;-) at all, and also emit diagnostics for OpenACC kernels constructs, and also added considerably more testsuite coverage. Committed to trunk in r267146: commit 75180da2a558d3106e26173326933f65b417182c Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> Date: Fri Dec 14 20:41:58 2018 +0000 Add user-friendly diagnostics for OpenACC loop parallelism assigned gcc/ * omp-offload.c (inform_oacc_loop): New function. (execute_oacc_device_lower): Use it to display loop parallelism. gcc/testsuite/ * c-c++-common/goacc/note-parallelism.c: New test. * gfortran.dg/goacc/note-parallelism.f90: New test. * 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/kernels-1.c: Likewise. * c-c++-common/goacc/kernels-double-reduction-n.c: Likewise. * c-c++-common/goacc/kernels-double-reduction.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/kernels-loop-inner.f95: Likewise. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@267146 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog | 6 + gcc/omp-offload.c | 49 +++++++- gcc/testsuite/ChangeLog | 18 +++ .../goacc/classify-kernels-unparallelized.c | 3 +- .../c-c++-common/goacc/classify-kernels.c | 3 +- .../c-c++-common/goacc/classify-parallel.c | 3 +- .../c-c++-common/goacc/classify-routine.c | 3 +- gcc/testsuite/c-c++-common/goacc/kernels-1.c | 10 +- .../goacc/kernels-double-reduction-n.c | 3 +- .../c-c++-common/goacc/kernels-double-reduction.c | 3 +- .../c-c++-common/goacc/note-parallelism.c | 115 ++++++++++++++++++ .../goacc/classify-kernels-unparallelized.f95 | 3 +- .../gfortran.dg/goacc/classify-kernels.f95 | 3 +- .../gfortran.dg/goacc/classify-parallel.f95 | 3 +- .../gfortran.dg/goacc/classify-routine.f95 | 3 +- .../gfortran.dg/goacc/kernels-loop-inner.f95 | 3 +- .../gfortran.dg/goacc/note-parallelism.f90 | 131 +++++++++++++++++++++ 17 files changed, 346 insertions(+), 16 deletions(-) diff --git gcc/ChangeLog gcc/ChangeLog index 527164c4f9ec..7fb4958da485 100644 --- gcc/ChangeLog +++ gcc/ChangeLog @@ -1,3 +1,9 @@ +2018-12-14 Thomas Schwinge <thomas@codesourcery.com> + Cesar Philippidis <cesar@codesourcery.com> + + * omp-offload.c (inform_oacc_loop): New function. + (execute_oacc_device_lower): Use it to display loop parallelism. + 2018-12-14 Jakub Jelinek <jakub@redhat.com> PR c++/82294 diff --git gcc/omp-offload.c gcc/omp-offload.c index 0abf0283c9e2..4457e1a3079b 100644 --- gcc/omp-offload.c +++ gcc/omp-offload.c @@ -823,7 +823,7 @@ dump_oacc_loop_part (FILE *file, gcall *from, int depth, } } -/* Dump OpenACC loops LOOP, its siblings and its children. */ +/* Dump OpenACC loop LOOP, its children, and its siblings. */ static void dump_oacc_loop (FILE *file, oacc_loop *loop, int depth) @@ -866,6 +866,31 @@ debug_oacc_loop (oacc_loop *loop) dump_oacc_loop (stderr, loop, 0); } +/* Provide diagnostics on OpenACC loop LOOP, its children, and its + siblings. */ + +static void +inform_oacc_loop (const oacc_loop *loop) +{ + const char *gang + = loop->mask & GOMP_DIM_MASK (GOMP_DIM_GANG) ? " gang" : ""; + const char *worker + = loop->mask & GOMP_DIM_MASK (GOMP_DIM_WORKER) ? " worker" : ""; + const char *vector + = loop->mask & GOMP_DIM_MASK (GOMP_DIM_VECTOR) ? " vector" : ""; + const char *seq = loop->mask == 0 ? " seq" : ""; + const dump_user_location_t loc + = dump_user_location_t::from_location_t (loop->loc); + dump_printf_loc (MSG_OPTIMIZED_LOCATIONS, loc, + "assigned OpenACC%s%s%s%s loop parallelism\n", gang, worker, + vector, seq); + + if (loop->child) + inform_oacc_loop (loop->child); + if (loop->sibling) + inform_oacc_loop (loop->sibling); +} + /* DFS walk of basic blocks BB onwards, creating OpenACC loop structures as we go. By construction these loops are properly nested. */ @@ -1533,6 +1558,28 @@ execute_oacc_device_lower () dump_oacc_loop (dump_file, loops, 0); fprintf (dump_file, "\n"); } + if (dump_enabled_p ()) + { + oacc_loop *l = loops; + /* OpenACC kernels constructs are special: they currently don't use the + generic oacc_loop infrastructure. */ + if (is_oacc_kernels) + { + /* Create a fake oacc_loop for diagnostic purposes. */ + l = new_oacc_loop_raw (NULL, + DECL_SOURCE_LOCATION (current_function_decl)); + l->mask = used_mask; + } + else + { + /* Skip the outermost, dummy OpenACC loop */ + l = l->child; + } + if (l) + inform_oacc_loop (l); + if (is_oacc_kernels) + free_oacc_loop (l); + } /* Offloaded targets may introduce new basic blocks, which require dominance information to update SSA. */ diff --git gcc/testsuite/ChangeLog gcc/testsuite/ChangeLog index 4af1cf9423d2..6b035b34ecfd 100644 --- gcc/testsuite/ChangeLog +++ gcc/testsuite/ChangeLog @@ -1,3 +1,21 @@ +2018-12-14 Thomas Schwinge <thomas@codesourcery.com> + Cesar Philippidis <cesar@codesourcery.com> + + * c-c++-common/goacc/note-parallelism.c: New test. + * gfortran.dg/goacc/note-parallelism.f90: New test. + * 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/kernels-1.c: Likewise. + * c-c++-common/goacc/kernels-double-reduction-n.c: Likewise. + * c-c++-common/goacc/kernels-double-reduction.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/kernels-loop-inner.f95: Likewise. + 2018-12-14 Alexandre Oliva <aoliva@redhat.com> PR c++/86823 diff --git gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c index 626f6b4fe687..d4c4b2ca237a 100644 --- gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c +++ gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c @@ -2,6 +2,7 @@ OpenACC kernels. */ /* { dg-additional-options "-O2" } + { dg-additional-options "-fopt-info-optimized-omp" } { dg-additional-options "-fdump-tree-ompexp" } { dg-additional-options "-fdump-tree-parloops1-all" } { dg-additional-options "-fdump-tree-oaccdevlow" } */ @@ -18,7 +19,7 @@ extern unsigned int f (unsigned int); void KERNELS () { -#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) +#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++) c[i] = a[f (i)] + b[f (i)]; } diff --git gcc/testsuite/c-c++-common/goacc/classify-kernels.c gcc/testsuite/c-c++-common/goacc/classify-kernels.c index 95037e6120ea..16e9b9e31d16 100644 --- gcc/testsuite/c-c++-common/goacc/classify-kernels.c +++ gcc/testsuite/c-c++-common/goacc/classify-kernels.c @@ -2,6 +2,7 @@ kernels. */ /* { dg-additional-options "-O2" } + { dg-additional-options "-fopt-info-optimized-omp" } { dg-additional-options "-fdump-tree-ompexp" } { dg-additional-options "-fdump-tree-parloops1-all" } { dg-additional-options "-fdump-tree-oaccdevlow" } */ @@ -14,7 +15,7 @@ extern unsigned int *__restrict c; void KERNELS () { -#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) +#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ for (unsigned int i = 0; i < N; i++) c[i] = a[i] + b[i]; } diff --git gcc/testsuite/c-c++-common/goacc/classify-parallel.c gcc/testsuite/c-c++-common/goacc/classify-parallel.c index 4f97301f5cf0..66a6d1336638 100644 --- gcc/testsuite/c-c++-common/goacc/classify-parallel.c +++ gcc/testsuite/c-c++-common/goacc/classify-parallel.c @@ -2,6 +2,7 @@ parallel. */ /* { dg-additional-options "-O2" } + { dg-additional-options "-fopt-info-optimized-omp" } { dg-additional-options "-fdump-tree-ompexp" } { dg-additional-options "-fdump-tree-oaccdevlow" } */ @@ -13,7 +14,7 @@ extern unsigned int *__restrict c; void PARALLEL () { -#pragma acc parallel loop copyin (a[0:N], b[0:N]) copyout (c[0:N]) +#pragma acc parallel loop copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */ for (unsigned int i = 0; i < N; i++) c[i] = a[i] + b[i]; } diff --git gcc/testsuite/c-c++-common/goacc/classify-routine.c gcc/testsuite/c-c++-common/goacc/classify-routine.c index fd89fc1ec662..a723d2cdf513 100644 --- gcc/testsuite/c-c++-common/goacc/classify-routine.c +++ gcc/testsuite/c-c++-common/goacc/classify-routine.c @@ -2,6 +2,7 @@ routine. */ /* { dg-additional-options "-O2" } + { dg-additional-options "-fopt-info-optimized-omp" } { dg-additional-options "-fdump-tree-ompexp" } { dg-additional-options "-fdump-tree-oaccdevlow" } */ @@ -15,7 +16,7 @@ extern unsigned int *__restrict c; #pragma acc routine worker void ROUTINE () { -#pragma acc loop +#pragma acc loop /* { dg-message "optimized: assigned OpenACC worker vector loop parallelism" } */ for (unsigned int i = 0; i < N; i++) c[i] = a[i] + b[i]; } diff --git gcc/testsuite/c-c++-common/goacc/kernels-1.c gcc/testsuite/c-c++-common/goacc/kernels-1.c index 4fcf86eca698..016abbdfe8d8 100644 --- gcc/testsuite/c-c++-common/goacc/kernels-1.c +++ gcc/testsuite/c-c++-common/goacc/kernels-1.c @@ -1,7 +1,9 @@ +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + int kernels_empty (void) { -#pragma acc kernels +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ ; return 0; @@ -10,7 +12,7 @@ kernels_empty (void) int kernels_eternal (void) { -#pragma acc kernels +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ { while (1) ; @@ -22,7 +24,7 @@ kernels_eternal (void) int kernels_noreturn (void) { -#pragma acc kernels +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ __builtin_abort (); return 0; @@ -36,7 +38,7 @@ kernels_loop_ptr_it (void) { float *i; -#pragma acc kernels +#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ { #pragma acc loop for (i = &b[0][0][0]; i < &b[0][0][10]; i++) diff --git gcc/testsuite/c-c++-common/goacc/kernels-double-reduction-n.c gcc/testsuite/c-c++-common/goacc/kernels-double-reduction-n.c index 10b364b367c0..8f7f415b58d8 100644 --- gcc/testsuite/c-c++-common/goacc/kernels-double-reduction-n.c +++ gcc/testsuite/c-c++-common/goacc/kernels-double-reduction-n.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fopt-info-optimized-omp" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ @@ -14,7 +15,7 @@ foo (unsigned int n) int i, j; unsigned int sum = 1; -#pragma acc kernels copyin (a[0:n]) copy (sum) +#pragma acc kernels copyin (a[0:n]) copy (sum) /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ { for (i = 0; i < n; ++i) for (j = 0; j < n; ++j) diff --git gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c index c0263460c1dc..c11d36fb4373 100644 --- gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c +++ gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fopt-info-optimized-omp" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ @@ -14,7 +15,7 @@ foo (void) int i, j; unsigned int sum = 1; -#pragma acc kernels copyin (a[0:N]) copy (sum) +#pragma acc kernels copyin (a[0:N]) copy (sum) /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ { for (i = 0; i < N; ++i) for (j = 0; j < N; ++j) diff --git gcc/testsuite/c-c++-common/goacc/note-parallelism.c gcc/testsuite/c-c++-common/goacc/note-parallelism.c new file mode 100644 index 000000000000..735df7dfad7a --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/note-parallelism.c @@ -0,0 +1,115 @@ +/* Test the output of "-fopt-info-optimized-omp". */ + +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ + +int +main () +{ + int x, y, z; + +#pragma acc parallel + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop gang /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop worker /* { dg-message "optimized: assigned OpenACC worker loop parallelism" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop vector /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop gang vector /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop gang worker /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop worker vector /* { dg-message "optimized: assigned OpenACC worker vector loop parallelism" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop gang worker vector /* { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop gang /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ + for (x = 0; x < 10; x++) +#pragma acc loop worker /* { dg-message "optimized: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop vector /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc parallel loop /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + for (x = 0; x < 10; x++) +#pragma acc loop /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) + ; + +#pragma acc parallel loop /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */ + for (x = 0; x < 10; x++) +#pragma acc loop /* { dg-message "optimized: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc parallel + for (x = 0; x < 10; x++) +#pragma acc loop /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc parallel loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) +#pragma acc loop /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc parallel loop /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + for (x = 0; x < 10; x++) +#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc parallel loop /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */ + for (x = 0; x < 10; x++) +#pragma acc loop /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc parallel loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) +#pragma acc loop /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + + return 0; +} diff --git gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 index 4b282cab5ae4..08772428c4c5 100644 --- gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 +++ gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 @@ -2,6 +2,7 @@ ! OpenACC kernels. ! { dg-additional-options "-O2" } +! { dg-additional-options "-fopt-info-optimized-omp" } ! { dg-additional-options "-fdump-tree-ompexp" } ! { dg-additional-options "-fdump-tree-parloops1-all" } ! { dg-additional-options "-fdump-tree-oaccdevlow" } @@ -19,7 +20,7 @@ program main call setup(a, b) !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) - do i = 0, n - 1 + do i = 0, n - 1 ! { dg-message "optimized: assigned OpenACC seq loop parallelism" } c(i) = a(f (i)) + b(f (i)) end do !$acc end kernels diff --git gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 index da025c19af87..f2c4736e111c 100644 --- gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 +++ gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 @@ -2,6 +2,7 @@ ! kernels. ! { dg-additional-options "-O2" } +! { dg-additional-options "-fopt-info-optimized-omp" } ! { dg-additional-options "-fdump-tree-ompexp" } ! { dg-additional-options "-fdump-tree-parloops1-all" } ! { dg-additional-options "-fdump-tree-oaccdevlow" } @@ -15,7 +16,7 @@ program main call setup(a, b) !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) - do i = 0, n - 1 + do i = 0, n - 1 ! { dg-message "optimized: assigned OpenACC gang loop parallelism" } c(i) = a(i) + b(i) end do !$acc end kernels diff --git gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 index e215c79be618..a23ea81609b4 100644 --- gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 +++ gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 @@ -2,6 +2,7 @@ ! parallel. ! { dg-additional-options "-O2" } +! { dg-additional-options "-fopt-info-optimized-omp" } ! { dg-additional-options "-fdump-tree-ompexp" } ! { dg-additional-options "-fdump-tree-oaccdevlow" } @@ -13,7 +14,7 @@ program main call setup(a, b) - !$acc parallel loop copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) + !$acc parallel loop copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } do i = 0, n - 1 c(i) = a(i) + b(i) end do diff --git gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 index 4ca406742f90..5cf4c13acb81 100644 --- gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 +++ gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 @@ -2,6 +2,7 @@ ! routine. ! { dg-additional-options "-O2" } +! { dg-additional-options "-fopt-info-optimized-omp" } ! { dg-additional-options "-fdump-tree-ompexp" } ! { dg-additional-options "-fdump-tree-oaccdevlow" } @@ -13,7 +14,7 @@ subroutine ROUTINE call setup(a, b) - !$acc loop + !$acc loop ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" } do i = 0, n - 1 c(i) = a(i) + b(i) end do diff --git gcc/testsuite/gfortran.dg/goacc/kernels-loop-inner.f95 gcc/testsuite/gfortran.dg/goacc/kernels-loop-inner.f95 index 333474141526..a3ad591f926c 100644 --- gcc/testsuite/gfortran.dg/goacc/kernels-loop-inner.f95 +++ gcc/testsuite/gfortran.dg/goacc/kernels-loop-inner.f95 @@ -1,4 +1,5 @@ ! { dg-additional-options "-O2" } +! { dg-additional-options "-fopt-info-optimized-omp" } program main implicit none @@ -6,7 +7,7 @@ program main integer :: a(100,100), b(100,100) integer :: i, j, d - !$acc kernels + !$acc kernels ! { dg-message "optimized: assigned OpenACC seq loop parallelism" } do i=1,100 do j=1,100 a(i,j) = 1 diff --git gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 new file mode 100644 index 000000000000..6c8a1bdc6a4c --- /dev/null +++ gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 @@ -0,0 +1,131 @@ +! Test the output of "-fopt-info-optimized-omp". + +! { dg-additional-options "-fopt-info-optimized-omp" } + +! See also "../../c-c++-common/goacc/note-parallelism.c". + +program test + implicit none + + integer x, y, z + + !$acc parallel + do x = 1, 10 + end do + !$acc end parallel + + !$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" } + do x = 1, 10 + end do + + !$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelis" } + do x = 1, 10 + end do + + !$acc parallel loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" } + do x = 1, 10 + end do + + !$acc parallel loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" } + do x = 1, 10 + end do + + !$acc parallel loop gang vector ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } + do x = 1, 10 + end do + + !$acc parallel loop gang worker ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } + do x = 1, 10 + end do + + !$acc parallel loop worker vector ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" } + do x = 1, 10 + end do + + !$acc parallel loop gang worker vector ! { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" } + do x = 1, 10 + end do + + !$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" } + do x = 1, 10 + !$acc loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" } + do y = 1, 10 + !$acc loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" } + do z = 1, 10 + end do + end do + end do + + !$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } + do x = 1, 10 + end do + + !$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } + do x = 1, 10 + !$acc loop ! { dg-message "optimized: assigned OpenACC vector loop parallelism" } + do y = 1, 10 + end do + end do + + !$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang loop parallelism" } + do x = 1, 10 + !$acc loop ! { dg-message "optimized: assigned OpenACC worker loop parallelism" } + do y = 1, 10 + !$acc loop ! { dg-message "optimized: assigned OpenACC vector loop parallelism" } + do z = 1, 10 + end do + end do + end do + + !$acc parallel + do x = 1, 10 + !$acc loop ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } + do y = 1, 10 + !$acc loop ! { dg-message "optimized: assigned OpenACC vector loop parallelism" } + do z = 1, 10 + end do + end do + end do + !$acc end parallel + + !$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" } + do x = 1, 10 + !$acc loop ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } + do y = 1, 10 + !$acc loop ! { dg-message "optimized: assigned OpenACC vector loop parallelism" } + do z = 1, 10 + end do + end do + end do + + !$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } + do x = 1, 10 + !$acc loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" } + do y = 1, 10 + !$acc loop ! { dg-message "optimized: assigned OpenACC vector loop parallelism" } + do z = 1, 10 + end do + end do + end do + + !$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } + do x = 1, 10 + !$acc loop ! { dg-message "optimized: assigned OpenACC vector loop parallelism" } + do y = 1, 10 + !$acc loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" } + do z = 1, 10 + end do + end do + end do + + !$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" } + do x = 1, 10 + !$acc loop ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } + do y = 1, 10 + !$acc loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" } + do z = 1, 10 + end do + end do + end do + +end program test Grüße Thomas ^ permalink raw reply [flat|nested] 12+ messages in thread
* [og8] Add user-friendly diagnostics for OpenACC loop parallelism assigned 2018-12-14 21:03 ` Add user-friendly diagnostics for OpenACC loop parallelism assigned (was: [PATCH 3/3] Add user-friendly OpenACC diagnostics regarding detected parallelism) Thomas Schwinge @ 2019-01-28 21:51 ` Thomas Schwinge 0 siblings, 0 replies; 12+ messages in thread From: Thomas Schwinge @ 2019-01-28 21:51 UTC (permalink / raw) To: gcc-patches; +Cc: Gergö Barany [-- Attachment #1: Type: text/plain, Size: 703 bytes --] Hi! On Fri, 14 Dec 2018 22:03:26 +0100, I wrote: > On Thu, 26 Jul 2018 07:14:21 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote: > > [...] > > I further changed that to make it build ;-) at all, and also emit > diagnostics for OpenACC kernels constructs, and also added considerably > more testsuite coverage. > > Committed to trunk in r267146: > > commit 75180da2a558d3106e26173326933f65b417182c > Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> > Date: Fri Dec 14 20:41:58 2018 +0000 > > Add user-friendly diagnostics for OpenACC loop parallelism assigned Backported my changes to openacc-gcc-8-branch, see attached. Grüße Thomas [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Add-user-friendly-diagnostics-for-OpenACC-loop-paral.patch --] [-- Type: text/x-diff, Size: 56753 bytes --] From 010a7a9220c55c6bdf803617c72ed912a790267c Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Wed, 23 Jan 2019 15:22:58 +0100 Subject: [PATCH] Add user-friendly diagnostics for OpenACC loop parallelism assigned Backport changes from trunk r267146. gcc/ * omp-offload.c (inform_oacc_loop, execute_oacc_device_lower): Update. gcc/testsuite/ * 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/kernels-1.c: Likewise. * c-c++-common/goacc/kernels-double-reduction-n.c: Likewise. * c-c++-common/goacc/kernels-double-reduction.c: Likewise. * c-c++-common/goacc/loop-auto-3.c: Likewise. * c-c++-common/goacc/note-parallelism.c: Likewise. * c-c++-common/goacc/orphan-reductions-2.c: Likewise. * gcc.target/nvptx/oacc-autopar.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/kernels-loop-inner.f95: Likewise. * gfortran.dg/goacc/loop-auto-1.f90: Likewise. * gfortran.dg/goacc/note-parallelism.f90: Likewise. * gfortran.dg/goacc/orphan-reductions-2.f90: Likewise. --- gcc/ChangeLog.openacc | 5 + gcc/omp-offload.c | 51 ++++++--- gcc/testsuite/ChangeLog.openacc | 22 ++++ .../goacc/classify-kernels-unparallelized.c | 3 +- .../c-c++-common/goacc/classify-kernels.c | 3 +- .../c-c++-common/goacc/classify-parallel.c | 3 +- .../c-c++-common/goacc/classify-routine.c | 3 +- gcc/testsuite/c-c++-common/goacc/kernels-1.c | 10 +- .../goacc/kernels-double-reduction-n.c | 3 +- .../goacc/kernels-double-reduction.c | 3 +- .../c-c++-common/goacc/loop-auto-3.c | 48 ++++----- .../c-c++-common/goacc/note-parallelism.c | 86 ++++++++++++--- .../c-c++-common/goacc/orphan-reductions-2.c | 26 ++--- gcc/testsuite/gcc.target/nvptx/oacc-autopar.c | 54 +++++----- .../goacc/classify-kernels-unparallelized.f95 | 3 +- .../gfortran.dg/goacc/classify-kernels.f95 | 3 +- .../gfortran.dg/goacc/classify-parallel.f95 | 3 +- .../gfortran.dg/goacc/classify-routine.f95 | 3 +- .../gfortran.dg/goacc/kernels-loop-inner.f95 | 3 +- .../gfortran.dg/goacc/loop-auto-1.f90 | 48 ++++----- .../gfortran.dg/goacc/note-parallelism.f90 | 101 +++++++++++++++--- .../gfortran.dg/goacc/orphan-reductions-2.f90 | 26 ++--- 22 files changed, 346 insertions(+), 164 deletions(-) diff --git a/gcc/ChangeLog.openacc b/gcc/ChangeLog.openacc index 22cdb5ba753..28d60da45b2 100644 --- a/gcc/ChangeLog.openacc +++ b/gcc/ChangeLog.openacc @@ -1,3 +1,8 @@ +2019-01-28 Thomas Schwinge <thomas@codesourcery.com> + + * omp-offload.c (inform_oacc_loop, execute_oacc_device_lower): + Update. + 2019-01-09 Julian Brown <julian@codesourcery.com> * doc/invoke.texi: Update mention of OpenACC version to 2.6. diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c index d428c6fdb62..60a85cba613 100644 --- a/gcc/omp-offload.c +++ b/gcc/omp-offload.c @@ -873,7 +873,7 @@ dump_oacc_loop_part (FILE *file, gcall *from, int depth, } } -/* Dump OpenACC loops LOOP, its siblings and its children. */ +/* Dump OpenACC loop LOOP, its children, and its siblings. */ static void dump_oacc_loop (FILE *file, oacc_loop *loop, int depth) @@ -916,23 +916,22 @@ debug_oacc_loop (oacc_loop *loop) dump_oacc_loop (stderr, loop, 0); } -/* Provide diagnostics on OpenACC loops LOOP, its siblings and its - children. */ +/* Provide diagnostics on OpenACC loop LOOP, its children, and its + siblings. */ static void -inform_oacc_loop (oacc_loop *loop) +inform_oacc_loop (const oacc_loop *loop) { + const char *gang + = loop->mask & GOMP_DIM_MASK (GOMP_DIM_GANG) ? " gang" : ""; + const char *worker + = loop->mask & GOMP_DIM_MASK (GOMP_DIM_WORKER) ? " worker" : ""; + const char *vector + = loop->mask & GOMP_DIM_MASK (GOMP_DIM_VECTOR) ? " vector" : ""; const char *seq = loop->mask == 0 ? " seq" : ""; - const char *gang = loop->mask & GOMP_DIM_MASK (GOMP_DIM_GANG) - ? " gang" : ""; - const char *worker = loop->mask & GOMP_DIM_MASK (GOMP_DIM_WORKER) - ? " worker" : ""; - const char *vector = loop->mask & GOMP_DIM_MASK (GOMP_DIM_VECTOR) - ? " vector" : ""; - - dump_printf_loc (MSG_NOTE, loop->loc, - "Detected parallelism <acc loop%s%s%s%s>\n", seq, gang, - worker, vector); + dump_printf_loc (MSG_OPTIMIZED_LOCATIONS, loop->loc, + "assigned OpenACC%s%s%s%s loop parallelism\n", gang, worker, + vector, seq); if (loop->child) inform_oacc_loop (loop->child); @@ -1664,8 +1663,28 @@ execute_oacc_device_lower () dump_oacc_loop (dump_file, loops, 0); fprintf (dump_file, "\n"); } - if (dump_enabled_p () && loops->child) - inform_oacc_loop (loops->child); + if (dump_enabled_p ()) + { + oacc_loop *l = loops; + /* OpenACC kernels constructs are special: they currently don't use the + generic oacc_loop infrastructure. */ + if (is_oacc_kernels) + { + /* Create a fake oacc_loop for diagnostic purposes. */ + l = new_oacc_loop_raw (NULL, + DECL_SOURCE_LOCATION (current_function_decl)); + l->mask = used_mask; + } + else + { + /* Skip the outermost, dummy OpenACC loop */ + l = l->child; + } + if (l) + inform_oacc_loop (l); + if (is_oacc_kernels) + free_oacc_loop (l); + } /* Offloaded targets may introduce new basic blocks, which require dominance information to update SSA. */ diff --git a/gcc/testsuite/ChangeLog.openacc b/gcc/testsuite/ChangeLog.openacc index 3bdce2e457d..ea136a306dc 100644 --- a/gcc/testsuite/ChangeLog.openacc +++ b/gcc/testsuite/ChangeLog.openacc @@ -1,3 +1,25 @@ +2019-01-28 Thomas Schwinge <thomas@codesourcery.com> + + * 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/kernels-1.c: Likewise. + * c-c++-common/goacc/kernels-double-reduction-n.c: Likewise. + * c-c++-common/goacc/kernels-double-reduction.c: Likewise. + * c-c++-common/goacc/loop-auto-3.c: Likewise. + * c-c++-common/goacc/note-parallelism.c: Likewise. + * c-c++-common/goacc/orphan-reductions-2.c: Likewise. + * gcc.target/nvptx/oacc-autopar.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/kernels-loop-inner.f95: Likewise. + * gfortran.dg/goacc/loop-auto-1.f90: Likewise. + * gfortran.dg/goacc/note-parallelism.f90: Likewise. + * gfortran.dg/goacc/orphan-reductions-2.f90: Likewise. + 2019-01-09 Julian Brown <julian@codesourcery.com> * c-c++-common/cpp/openacc-define-3.c: Update expected value for 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 626f6b4fe68..64467774037 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c @@ -2,6 +2,7 @@ OpenACC kernels. */ /* { dg-additional-options "-O2" } + { dg-additional-options "-fopt-info-optimized-omp" } { dg-additional-options "-fdump-tree-ompexp" } { dg-additional-options "-fdump-tree-parloops1-all" } { dg-additional-options "-fdump-tree-oaccdevlow" } */ @@ -18,7 +19,7 @@ extern unsigned int f (unsigned int); void KERNELS () { -#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) +#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ for (unsigned int i = 0; i < N; i++) 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 95037e6120e..c59a65e1d0f 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c @@ -2,6 +2,7 @@ kernels. */ /* { dg-additional-options "-O2" } + { dg-additional-options "-fopt-info-optimized-omp" } { dg-additional-options "-fdump-tree-ompexp" } { dg-additional-options "-fdump-tree-parloops1-all" } { dg-additional-options "-fdump-tree-oaccdevlow" } */ @@ -14,7 +15,7 @@ extern unsigned int *__restrict c; void KERNELS () { -#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) +#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ for (unsigned int i = 0; i < N; i++) c[i] = a[i] + b[i]; } diff --git a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c index 4f97301f5cf..b345c225aea 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c @@ -2,6 +2,7 @@ parallel. */ /* { dg-additional-options "-O2" } + { dg-additional-options "-fopt-info-optimized-omp" } { dg-additional-options "-fdump-tree-ompexp" } { dg-additional-options "-fdump-tree-oaccdevlow" } */ @@ -13,7 +14,7 @@ extern unsigned int *__restrict c; void PARALLEL () { -#pragma acc parallel loop copyin (a[0:N], b[0:N]) copyout (c[0:N]) +#pragma acc parallel loop copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ for (unsigned int i = 0; i < N; i++) c[i] = a[i] + b[i]; } diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine.c b/gcc/testsuite/c-c++-common/goacc/classify-routine.c index f54c3942bbf..5ca2ec9c603 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-routine.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-routine.c @@ -2,6 +2,7 @@ routine. */ /* { dg-additional-options "-O2" } + { dg-additional-options "-fopt-info-optimized-omp" } { dg-additional-options "-fdump-tree-ompexp" } { dg-additional-options "-fdump-tree-oaccdevlow" } */ @@ -15,7 +16,7 @@ extern unsigned int *__restrict c; #pragma acc routine worker void ROUTINE () { -#pragma acc loop +#pragma acc loop /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */ for (unsigned int i = 0; i < N; i++) c[i] = a[i] + b[i]; } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-1.c index 4fcf86eca69..0a4bd854611 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-1.c @@ -1,7 +1,9 @@ +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + int kernels_empty (void) { -#pragma acc kernels +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ ; return 0; @@ -10,7 +12,7 @@ kernels_empty (void) int kernels_eternal (void) { -#pragma acc kernels +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ { while (1) ; @@ -22,7 +24,7 @@ kernels_eternal (void) int kernels_noreturn (void) { -#pragma acc kernels +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ __builtin_abort (); return 0; @@ -36,7 +38,7 @@ kernels_loop_ptr_it (void) { float *i; -#pragma acc kernels +#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ { #pragma acc loop for (i = &b[0][0][0]; i < &b[0][0][10]; i++) diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-double-reduction-n.c b/gcc/testsuite/c-c++-common/goacc/kernels-double-reduction-n.c index 10b364b367c..dd3b7c8b144 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-double-reduction-n.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-double-reduction-n.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fopt-info-optimized-omp" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ @@ -14,7 +15,7 @@ foo (unsigned int n) int i, j; unsigned int sum = 1; -#pragma acc kernels copyin (a[0:n]) copy (sum) +#pragma acc kernels copyin (a[0:n]) copy (sum) /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ { for (i = 0; i < n; ++i) for (j = 0; j < n; ++j) diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c b/gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c index c0263460c1d..0175434a20c 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-double-reduction.c @@ -1,4 +1,5 @@ /* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fopt-info-optimized-omp" } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ @@ -14,7 +15,7 @@ foo (void) int i, j; unsigned int sum = 1; -#pragma acc kernels copyin (a[0:N]) copy (sum) +#pragma acc kernels copyin (a[0:N]) copy (sum) /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ { for (i = 0; i < N; ++i) for (j = 0; j < N; ++j) diff --git a/gcc/testsuite/c-c++-common/goacc/loop-auto-3.c b/gcc/testsuite/c-c++-common/goacc/loop-auto-3.c index 2c32dcadbc0..97b8d75104c 100644 --- a/gcc/testsuite/c-c++-common/goacc/loop-auto-3.c +++ b/gcc/testsuite/c-c++-common/goacc/loop-auto-3.c @@ -1,78 +1,78 @@ /* Ensure that the auto clause falls back to seq parallelism when the OpenACC loop is not explicitly independent. */ -/* { dg-additional-options "-fopt-info-note-omp" } */ +/* { dg-additional-options "-fopt-info-optimized-omp" } */ void test () { int i, j, k, l, n = 100; -#pragma acc parallel loop auto /* { dg-message "Detected parallelism <acc loop seq>" } */ +#pragma acc parallel loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ for (i = 0; i < n; i++) -#pragma acc loop auto independent /* { dg-message "Detected parallelism <acc loop gang>" } */ +#pragma acc loop auto independent /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ for (j = 0; j < n; j++) -#pragma acc loop worker vector /* { dg-message "Detected parallelism <acc loop worker vector>" } */ +#pragma acc loop worker vector /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */ for (k = 0; k < n; k++) ; -#pragma acc parallel loop auto independent /* { dg-message "Detected parallelism <acc loop gang worker>" } */ +#pragma acc parallel loop auto independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ for (i = 0; i < n; i++) -#pragma acc loop auto /* { dg-message "Detected parallelism <acc loop seq>" } */ +#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ for (j = 0; j < n; j++) -#pragma acc loop auto /* { dg-message "Detected parallelism <acc loop seq>" } */ +#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ for (k = 0; k < n; k++) -#pragma acc loop auto independent /* { dg-message "Detected parallelism <acc loop vector>" } */ +#pragma acc loop auto independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (l = 0; l < n; l++) ; -#pragma acc parallel loop gang /* { dg-message "Detected parallelism <acc loop gang>" } */ +#pragma acc parallel loop gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ for (i = 0; i < n; i++) -#pragma acc loop worker /* { dg-message "Detected parallelism <acc loop worker>" } */ +#pragma acc loop worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ for (j = 0; j < n; j++) -#pragma acc loop vector /* { dg-message "Detected parallelism <acc loop vector>" } */ +#pragma acc loop vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (k = 0; k < n; k++) { -#pragma acc loop auto independent /* { dg-message "Detected parallelism <acc loop seq>" } */ +#pragma acc loop auto independent /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ /* { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } */ for (l = 0; l < n; l++) ; -#pragma acc loop auto /* { dg-message "Detected parallelism <acc loop seq>" } */ +#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ for (l = 0; l < n; l++) ; } -#pragma acc parallel loop /* { dg-message "Detected parallelism <acc loop seq>" } */ +#pragma acc parallel loop /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ /* { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } */ for (i = 0; i < n; i++) { -#pragma acc loop gang worker /* { dg-message "Detected parallelism <acc loop gang worker>" } */ +#pragma acc loop gang worker /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ for (j = 0; j < n; j++) -#pragma acc loop auto /* { dg-message "Detected parallelism <acc loop seq>" } */ +#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ for (k = 0; k < n; k++) { -#pragma acc loop vector /* { dg-message "Detected parallelism <acc loop vector>" } */ +#pragma acc loop vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (l = 0; l < n; l++) ; -#pragma acc loop auto independent /* { dg-message "Detected parallelism <acc loop vector>" } */ +#pragma acc loop auto independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (l = 0; l < n; l++) ; } -#pragma acc loop worker /* { dg-message "Detected parallelism <acc loop worker>" } */ +#pragma acc loop worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ for (j = 0; j < n; j++) -#pragma acc loop vector /* { dg-message "Detected parallelism <acc loop vector>" } */ +#pragma acc loop vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (k = 0; k < n; k++) ; } -#pragma acc parallel loop /* { dg-message "Detected parallelism <acc loop gang>" } */ +#pragma acc parallel loop /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ for (i = 0; i < n; i++) -#pragma acc loop /* { dg-message "Detected parallelism <acc loop worker>" } */ +#pragma acc loop /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ for (j = 0; j < n; j++) -#pragma acc loop /* { dg-message "Detected parallelism <acc loop seq>" } */ +#pragma acc loop /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ /* { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } */ for (k = 0; k < n; k++) -#pragma acc loop /* { dg-message "Detected parallelism <acc loop vector>" } */ +#pragma acc loop /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (l = 0; l < n; l++) ; } diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism.c b/gcc/testsuite/c-c++-common/goacc/note-parallelism.c index 3ec794c7ab4..7548fb72d14 100644 --- a/gcc/testsuite/c-c++-common/goacc/note-parallelism.c +++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism.c @@ -1,59 +1,113 @@ -/* Test the output of -fopt-info-note-omp. */ +/* Test the output of "-fopt-info-optimized-omp". */ -/* { dg-additional-options "-fopt-info-note-omp" } */ +/* { dg-additional-options "-fopt-info-optimized-omp" } */ + +/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */ int main () { int x, y, z; -#pragma acc parallel loop seq /* { dg-message "note: Detected parallelism <acc loop seq>" } */ +#pragma acc parallel for (x = 0; x < 10; x++) ; -#pragma acc parallel loop gang /* { dg-message "note: Detected parallelism <acc loop gang>" } */ +#pragma acc parallel loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ for (x = 0; x < 10; x++) ; -#pragma acc parallel loop worker /* { dg-message "note: Detected parallelism <acc loop worker>" } */ +#pragma acc parallel loop gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ for (x = 0; x < 10; x++) ; -#pragma acc parallel loop vector /* { dg-message "note: Detected parallelism <acc loop vector>" } */ +#pragma acc parallel loop worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ for (x = 0; x < 10; x++) ; -#pragma acc parallel loop gang vector /* { dg-message "note: Detected parallelism <acc loop gang vector>" } */ +#pragma acc parallel loop vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (x = 0; x < 10; x++) ; -#pragma acc parallel loop gang worker /* { dg-message "note: Detected parallelism <acc loop gang worker>" } */ +#pragma acc parallel loop gang vector /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ for (x = 0; x < 10; x++) ; -#pragma acc parallel loop worker vector /* { dg-message "note: Detected parallelism <acc loop worker vector>" } */ +#pragma acc parallel loop gang worker /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ for (x = 0; x < 10; x++) ; -#pragma acc parallel loop gang worker vector /* { dg-message "note: Detected parallelism <acc loop gang worker vector>" } */ +#pragma acc parallel loop worker vector /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */ for (x = 0; x < 10; x++) ; -#pragma acc parallel loop /* { dg-message "note: Detected parallelism <acc loop gang vector>" } */ +#pragma acc parallel loop gang worker vector /* { dg-message "note: assigned OpenACC gang worker vector loop parallelism" } */ + for (x = 0; x < 10; x++) + ; + +#pragma acc parallel loop gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ + for (x = 0; x < 10; x++) +#pragma acc loop worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc parallel loop /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ for (x = 0; x < 10; x++) ; -#pragma acc parallel loop /* { dg-message "note: Detected parallelism <acc loop gang worker>" } */ +#pragma acc parallel loop /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ for (x = 0; x < 10; x++) -#pragma acc loop /* { dg-message "note: Detected parallelism <acc loop vector>" } */ +#pragma acc loop /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (y = 0; y < 10; y++) ; -#pragma acc parallel loop gang /* { dg-message "note: Detected parallelism <acc loop gang>" } */ +#pragma acc parallel loop /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ + for (x = 0; x < 10; x++) +#pragma acc loop /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc parallel + for (x = 0; x < 10; x++) +#pragma acc loop /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc parallel loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ + for (x = 0; x < 10; x++) +#pragma acc loop /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + for (y = 0; y < 10; y++) +#pragma acc loop /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc parallel loop /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + 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 /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ + for (z = 0; z < 10; z++) + ; + +#pragma acc parallel loop /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ + for (x = 0; x < 10; x++) +#pragma acc loop /* { 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 parallel loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ for (x = 0; x < 10; x++) -#pragma acc loop worker /* { dg-message "note: Detected parallelism <acc loop worker>" } */ +#pragma acc loop /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ for (y = 0; y < 10; y++) -#pragma acc loop vector /* { dg-message "note: Detected parallelism <acc loop vector>" } */ +#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ for (z = 0; z < 10; z++) ; diff --git a/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c index 8ea4ce37344..bc29f3672c4 100644 --- a/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c +++ b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c @@ -1,7 +1,7 @@ /* Ensure that the middle end does not assign gang level parallelism to orphan loop containing reductions. */ -/* { dg-additional-options "-fopt-info-note-omp" } */ +/* { dg-additional-options "-fopt-info-optimized-omp" } */ #pragma acc routine gang int @@ -9,7 +9,7 @@ f1 () /* { dg-warning "region is gang partitioned but does not contain gang part { int sum = 0, i; -#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <acc loop worker vector>" } */ +#pragma acc loop reduction (+:sum) /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */ for (i = 0; i < 100; i++) sum++; @@ -22,9 +22,9 @@ f2 () /* { dg-warning "region is gang partitioned but does not contain gang part { int sum = 0, i, j; -#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <acc loop worker>" } */ +#pragma acc loop reduction (+:sum) /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ for (i = 0; i < 100; i++) -#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <acc loop vector>" } */ +#pragma acc loop reduction (+:sum) /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (j = 0; j < 100; j++) sum++; @@ -37,12 +37,12 @@ f3 () /* { dg-warning "region is gang partitioned but does not contain gang part { int sum = 0, i, j, k; -#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <acc loop worker>" } */ +#pragma acc loop reduction (+:sum) /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ for (i = 0; i < 100; i++) -#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <acc loop seq>" } */ +#pragma acc loop reduction (+:sum) /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ /* { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } */ for (j = 0; j < 100; j++) -#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <acc loop vector>" } */ +#pragma acc loop reduction (+:sum) /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (k = 0; k < 100; k++) sum++; @@ -56,27 +56,27 @@ main () #pragma acc parallel copy (sum) { -#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <acc loop gang vector>" } */ +#pragma acc loop reduction (+:sum) /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ for (i = 0; i < 100; i++) sum++; } #pragma acc parallel copy (sum) { -#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <acc loop gang worker>" } */ +#pragma acc loop reduction (+:sum) /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ for (i = 0; i < 100; i++) -#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <acc loop vector>" } */ +#pragma acc loop reduction (+:sum) /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (j = 0; j < 100; j++) sum++; } #pragma acc parallel copy (sum) { -#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <acc loop gang>" } */ +#pragma acc loop reduction (+:sum) /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ for (i = 0; i < 100; i++) -#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <acc loop worker>" } */ +#pragma acc loop reduction (+:sum) /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ for (j = 0; j < 100; j++) -#pragma acc loop reduction (+:sum) /* { dg-message "Detected parallelism <acc loop vector>" } */ +#pragma acc loop reduction (+:sum) /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (k = 0; k < 100; k++) sum++; } diff --git a/gcc/testsuite/gcc.target/nvptx/oacc-autopar.c b/gcc/testsuite/gcc.target/nvptx/oacc-autopar.c index e574fb1838e..2cae7149eef 100644 --- a/gcc/testsuite/gcc.target/nvptx/oacc-autopar.c +++ b/gcc/testsuite/gcc.target/nvptx/oacc-autopar.c @@ -1,7 +1,7 @@ /* Verify the default assignment of parallelism. */ /* { dg-do compile } */ -/* { dg-additional-options "-fopenacc -fopt-info-note-omp" } */ +/* { dg-additional-options "-fopenacc -fopt-info-optimized-omp" } */ /* { dg-skip-if "cc1: error: option -mgomp is not supported together with -fopenacc" { *-*-* } { "-mgomp" } } */ void @@ -10,16 +10,16 @@ gang_independent () int i, j; #pragma acc parallel - #pragma acc loop gang /* { dg-message "Detected parallelism <acc loop gang>" } */ + #pragma acc loop gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ for (i = 0; i < 10; i++) - #pragma acc loop /* { dg-message "Detected parallelism <acc loop worker vector>" } */ + #pragma acc loop /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */ for (j = 0; j < 10; j++) ; #pragma acc parallel vector_length (128) - #pragma acc loop gang /* { dg-message "Detected parallelism <acc loop gang>" } */ + #pragma acc loop gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ for (i = 0; i < 10; i++) - #pragma acc loop /* { dg-message "Detected parallelism <acc loop vector>" } */ + #pragma acc loop /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (j = 0; j < 10; j++) ; } @@ -30,20 +30,20 @@ gang_independent_seq () int i, j, k; #pragma acc parallel - #pragma acc loop gang /* { dg-message "Detected parallelism <acc loop gang>" } */ + #pragma acc loop gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ for (i = 0; i < 10; i++) - #pragma acc loop /* { dg-message "Detected parallelism <acc loop worker vector>" } */ + #pragma acc loop /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */ for (j = 0; j < 10; j++) - #pragma acc loop seq /* { dg-message "Detected parallelism <acc loop seq>" } */ + #pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ for (k = 0; k < 10; k++) ; #pragma acc parallel vector_length (128) - #pragma acc loop gang /* { dg-message "Detected parallelism <acc loop gang>" } */ + #pragma acc loop gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ for (i = 0; i < 10; i++) - #pragma acc loop /* { dg-message "Detected parallelism <acc loop vector>" } */ + #pragma acc loop /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (j = 0; j < 10; j++) - #pragma acc loop seq /* { dg-message "Detected parallelism <acc loop seq>" } */ + #pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */ for (k = 0; k < 10; k++) ; } @@ -54,16 +54,16 @@ worker () int i, j; #pragma acc parallel - #pragma acc loop worker /* { dg-message "Detected parallelism <acc loop worker>" } */ + #pragma acc loop worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ for (i = 0; i < 10; i++) - #pragma acc loop /* { dg-message "Detected parallelism <acc loop vector>" } */ + #pragma acc loop /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (j = 0; j < 10; j++) ; #pragma acc parallel vector_length (128) - #pragma acc loop worker /* { dg-message "Detected parallelism <acc loop worker>" } */ + #pragma acc loop worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ for (i = 0; i < 10; i++) - #pragma acc loop /* { dg-message "Detected parallelism <acc loop vector>" } */ + #pragma acc loop /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (j = 0; j < 10; j++) ; } @@ -74,48 +74,48 @@ fully_independent () int i, j, k; #pragma acc parallel - #pragma acc loop /* { dg-message "Detected parallelism <acc loop gang vector>" } */ + #pragma acc loop /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ for (i = 0; i < 10; i++) ; #pragma acc parallel vector_length (128) - #pragma acc loop /* { dg-message "Detected parallelism <acc loop gang vector>" } */ + #pragma acc loop /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */ for (i = 0; i < 10; i++) ; #pragma acc parallel - #pragma acc loop /* { dg-message "Detected parallelism <acc loop gang worker>" } */ + #pragma acc loop /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ for (i = 0; i < 10; i++) - #pragma acc loop /* { dg-message "Detected parallelism <acc loop vector>" } */ + #pragma acc loop /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (j = 0; j < 10; j++) ; /* FIXME: Should the outer loop only be gang partitioned so that the inner loopp can utilize a large vector_length? */ #pragma acc parallel vector_length (128) - #pragma acc loop /* { dg-message "Detected parallelism <acc loop gang worker>" } */ + #pragma acc loop /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */ for (i = 0; i < 10; i++) - #pragma acc loop /* { dg-message "Detected parallelism <acc loop vector>" } */ + #pragma acc loop /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (j = 0; j < 10; j++) ; #pragma acc parallel - #pragma acc loop /* { dg-message "Detected parallelism <acc loop gang>" } */ + #pragma acc loop /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ for (i = 0; i < 10; i++) - #pragma acc loop /* { dg-message "Detected parallelism <acc loop worker>" } */ + #pragma acc loop /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ for (j = 0; j < 10; j++) - #pragma acc loop /* { dg-message "Detected parallelism <acc loop vector>" } */ + #pragma acc loop /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (k = 0; k < 10; k++) ; /* FIXME: Should the middle loop be seq-partitioned in order to respect vector_length = 128 on the innermost loop? */ #pragma acc parallel vector_length (128) - #pragma acc loop /* { dg-message "Detected parallelism <acc loop gang>" } */ + #pragma acc loop /* { dg-message "note: assigned OpenACC gang loop parallelism" } */ for (i = 0; i < 10; i++) - #pragma acc loop /* { dg-message "Detected parallelism <acc loop worker>" } */ + #pragma acc loop /* { dg-message "note: assigned OpenACC worker loop parallelism" } */ for (j = 0; j < 10; j++) - #pragma acc loop /* { dg-message "Detected parallelism <acc loop vector>" } */ + #pragma acc loop /* { dg-message "note: assigned OpenACC vector loop parallelism" } */ for (k = 0; k < 10; k++) ; } diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 index 4b282cab5ae..1b3b02aaf6c 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 @@ -2,6 +2,7 @@ ! OpenACC kernels. ! { dg-additional-options "-O2" } +! { dg-additional-options "-fopt-info-optimized-omp" } ! { dg-additional-options "-fdump-tree-ompexp" } ! { dg-additional-options "-fdump-tree-parloops1-all" } ! { dg-additional-options "-fdump-tree-oaccdevlow" } @@ -19,7 +20,7 @@ program main call setup(a, b) !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) - do i = 0, n - 1 + do i = 0, n - 1 ! { dg-message "note: assigned OpenACC seq loop parallelism" } c(i) = a(f (i)) + b(f (i)) end do !$acc end kernels diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 index da025c19af8..9d9827dbdf5 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 @@ -2,6 +2,7 @@ ! kernels. ! { dg-additional-options "-O2" } +! { dg-additional-options "-fopt-info-optimized-omp" } ! { dg-additional-options "-fdump-tree-ompexp" } ! { dg-additional-options "-fdump-tree-parloops1-all" } ! { dg-additional-options "-fdump-tree-oaccdevlow" } @@ -15,7 +16,7 @@ program main call setup(a, b) !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) - do i = 0, n - 1 + do i = 0, n - 1 ! { dg-message "note: assigned OpenACC gang loop parallelism" } c(i) = a(i) + b(i) end do !$acc end kernels diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 index e215c79be61..c1ed044b448 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 @@ -2,6 +2,7 @@ ! parallel. ! { dg-additional-options "-O2" } +! { dg-additional-options "-fopt-info-optimized-omp" } ! { dg-additional-options "-fdump-tree-ompexp" } ! { dg-additional-options "-fdump-tree-oaccdevlow" } @@ -13,7 +14,7 @@ program main call setup(a, b) - !$acc parallel loop copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) + !$acc parallel loop copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) ! { dg-message "note: assigned OpenACC gang vector loop parallelism" } do i = 0, n - 1 c(i) = a(i) + b(i) end do diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 index 445ff9afd4e..38c32fa4f65 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 @@ -2,6 +2,7 @@ ! routine. ! { dg-additional-options "-O2" } +! { dg-additional-options "-fopt-info-optimized-omp" } ! { dg-additional-options "-fdump-tree-ompexp" } ! { dg-additional-options "-fdump-tree-oaccdevlow" } @@ -13,7 +14,7 @@ subroutine ROUTINE call setup(a, b) - !$acc loop + !$acc loop ! { dg-message "note: assigned OpenACC worker vector loop parallelism" } do i = 0, n - 1 c(i) = a(i) + b(i) end do diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-inner.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-inner.f95 index 33347414152..16a64e6d76d 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-inner.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-inner.f95 @@ -1,4 +1,5 @@ ! { dg-additional-options "-O2" } +! { dg-additional-options "-fopt-info-optimized-omp" } program main implicit none @@ -6,7 +7,7 @@ program main integer :: a(100,100), b(100,100) integer :: i, j, d - !$acc kernels + !$acc kernels ! { dg-message "note: assigned OpenACC seq loop parallelism" } do i=1,100 do j=1,100 a(i,j) = 1 diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-auto-1.f90 b/gcc/testsuite/gfortran.dg/goacc/loop-auto-1.f90 index 7fc9210997c..2bbb6f79a05 100644 --- a/gcc/testsuite/gfortran.dg/goacc/loop-auto-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/loop-auto-1.f90 @@ -1,47 +1,47 @@ ! Ensure that the auto clause falls back to seq parallelism when the ! OpenACC loop is not explicitly independent. -! { dg-additional-options "-fopt-info-note-omp" } +! { dg-additional-options "-fopt-info-optimized-omp" } program test implicit none integer, parameter :: n = 100 integer i, j, k, l - !$acc parallel loop auto ! { dg-message "Detected parallelism <acc loop seq>" } + !$acc parallel loop auto ! { dg-message "note: assigned OpenACC seq loop parallelism" } do i = 1, n - !$acc loop auto independent ! { dg-message "Detected parallelism <acc loop gang>" } + !$acc loop auto independent ! { dg-message "note: assigned OpenACC gang loop parallelism" } do j = 1, n - !$acc loop worker vector ! { dg-message "Detected parallelism <acc loop worker vector>" } + !$acc loop worker vector ! { dg-message "note: assigned OpenACC worker vector loop parallelism" } do k = 1, n end do end do end do - !$acc parallel loop auto independent ! { dg-message "Detected parallelism <acc loop gang worker>" } + !$acc parallel loop auto independent ! { dg-message "note: assigned OpenACC gang worker loop parallelism" } do i = 1, n - !$acc loop auto ! { dg-message "Detected parallelism <acc loop seq>" } + !$acc loop auto ! { dg-message "note: assigned OpenACC seq loop parallelism" } do j = 1, n - !$acc loop auto ! { dg-message "Detected parallelism <acc loop seq>" } + !$acc loop auto ! { dg-message "note: assigned OpenACC seq loop parallelism" } do k = 1, n - !$acc loop auto independent ! { dg-message "Detected parallelism <acc loop vector>" } + !$acc loop auto independent ! { dg-message "note: assigned OpenACC vector loop parallelism" } do l = 1, n end do end do end do end do - !$acc parallel loop gang ! { dg-message "Detected parallelism <acc loop gang>" } + !$acc parallel loop gang ! { dg-message "note: assigned OpenACC gang loop parallelism" } do i = 1, n - !$acc loop worker ! { dg-message "Detected parallelism <acc loop worker>" } + !$acc loop worker ! { dg-message "note: assigned OpenACC worker loop parallelism" } do j = 1, n - !$acc loop vector ! { dg-message "Detected parallelism <acc loop vector>" } + !$acc loop vector ! { dg-message "note: assigned OpenACC vector loop parallelism" } do k = 1, n - !$acc loop auto independent ! { dg-message "Detected parallelism <acc loop seq>" } + !$acc loop auto independent ! { dg-message "note: assigned OpenACC seq loop parallelism" } ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } do l = 1, n end do - !$acc loop auto ! { dg-message "Detected parallelism <acc loop seq>" } + !$acc loop auto ! { dg-message "note: assigned OpenACC seq loop parallelism" } do l = 1, n end do end do @@ -49,37 +49,37 @@ program test end do - !$acc parallel loop ! { dg-message "Detected parallelism <acc loop seq>" } + !$acc parallel loop ! { dg-message "note: assigned OpenACC seq loop parallelism" } ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } do i = 1, n - !$acc loop gang worker ! { dg-message "Detected parallelism <acc loop gang worker>" } + !$acc loop gang worker ! { dg-message "note: assigned OpenACC gang worker loop parallelism" } do j = 1, n - !$acc loop auto ! { dg-message "Detected parallelism <acc loop seq>" } + !$acc loop auto ! { dg-message "note: assigned OpenACC seq loop parallelism" } do k = 1, n - !$acc loop vector ! { dg-message "Detected parallelism <acc loop vector>" } + !$acc loop vector ! { dg-message "note: assigned OpenACC vector loop parallelism" } do l = 1, n end do end do - !$acc loop auto independent ! { dg-message "Detected parallelism <acc loop vector>" } + !$acc loop auto independent ! { dg-message "note: assigned OpenACC vector loop parallelism" } do l = 1, n end do end do - !$acc loop worker ! { dg-message "Detected parallelism <acc loop worker>" } + !$acc loop worker ! { dg-message "note: assigned OpenACC worker loop parallelism" } do j = 1, n - !$acc loop vector ! { dg-message "Detected parallelism <acc loop vector>" } + !$acc loop vector ! { dg-message "note: assigned OpenACC vector loop parallelism" } do k = 1, n end do end do end do - !$acc parallel loop ! { dg-message "Detected parallelism <acc loop gang>" } + !$acc parallel loop ! { dg-message "note: assigned OpenACC gang loop parallelism" } do i = 1, n - !$acc loop ! { dg-message "Detected parallelism <acc loop worker>" } + !$acc loop ! { dg-message "note: assigned OpenACC worker loop parallelism" } do j = 1, n - !$acc loop ! { dg-message "Detected parallelism <acc loop seq>" } + !$acc loop ! { dg-message "note: assigned OpenACC seq loop parallelism" } ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } do k = 1, n - !$acc loop ! { dg-message "Detected parallelism <acc loop vector>" } + !$acc loop ! { dg-message "note: assigned OpenACC vector loop parallelism" } do l = 1, n end do end do diff --git a/gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 b/gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 index a0c78c52058..6aaf1fe290a 100644 --- a/gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/note-parallelism.f90 @@ -1,62 +1,131 @@ -! Test the output of -fopt-info-note-omp. +! Test the output of "-fopt-info-optimized-omp". -! { dg-additional-options "-fopt-info-note-omp" } +! { dg-additional-options "-fopt-info-optimized-omp" } + +! See also "../../c-c++-common/goacc/note-parallelism.c". program test implicit none integer x, y, z - !$acc parallel loop seq ! { dg-message "note: Detected parallelism <acc loop seq>" } + !$acc parallel + do x = 1, 10 + end do + !$acc end parallel + + !$acc parallel loop seq ! { dg-message "note: assigned OpenACC seq loop parallelism" } + do x = 1, 10 + end do + + !$acc parallel loop gang ! { dg-message "note: assigned OpenACC gang loop parallelis" } + do x = 1, 10 + end do + + !$acc parallel loop worker ! { dg-message "note: assigned OpenACC worker loop parallelism" } do x = 1, 10 end do - !$acc parallel loop gang ! { dg-message "note: Detected parallelism <acc loop gang>" } + !$acc parallel loop vector ! { dg-message "note: assigned OpenACC vector loop parallelism" } do x = 1, 10 end do - !$acc parallel loop worker ! { dg-message "note: Detected parallelism <acc loop worker>" } + !$acc parallel loop gang vector ! { dg-message "note: assigned OpenACC gang vector loop parallelism" } do x = 1, 10 end do - !$acc parallel loop vector ! { dg-message "note: Detected parallelism <acc loop vector>" } + !$acc parallel loop gang worker ! { dg-message "note: assigned OpenACC gang worker loop parallelism" } do x = 1, 10 end do - !$acc parallel loop gang vector ! { dg-message "note: Detected parallelism <acc loop gang vector>" } + !$acc parallel loop worker vector ! { dg-message "note: assigned OpenACC worker vector loop parallelism" } do x = 1, 10 end do - !$acc parallel loop gang worker ! { dg-message "note: Detected parallelism <acc loop gang worker>" } + !$acc parallel loop gang worker vector ! { dg-message "note: assigned OpenACC gang worker vector loop parallelism" } do x = 1, 10 end do - !$acc parallel loop worker vector ! { dg-message "note: Detected parallelism <acc loop worker vector>" } + !$acc parallel loop gang ! { dg-message "note: assigned OpenACC gang loop parallelism" } do x = 1, 10 + !$acc loop worker ! { dg-message "note: assigned OpenACC worker loop parallelism" } + do y = 1, 10 + !$acc loop vector ! { dg-message "note: assigned OpenACC vector loop parallelism" } + do z = 1, 10 + end do + end do end do - !$acc parallel loop gang worker vector ! { dg-message "note: Detected parallelism <acc loop gang worker vector>" } + !$acc parallel loop ! { dg-message "note: assigned OpenACC gang vector loop parallelism" } do x = 1, 10 end do - !$acc parallel loop ! { dg-message "note: Detected parallelism <acc loop gang vector>" } + !$acc parallel loop ! { dg-message "note: assigned OpenACC gang worker loop parallelism" } do x = 1, 10 + !$acc loop ! { dg-message "note: assigned OpenACC vector loop parallelism" } + do y = 1, 10 + end do end do - !$acc parallel loop ! { dg-message "note: Detected parallelism <acc loop gang worker>" } + !$acc parallel loop ! { dg-message "note: assigned OpenACC gang loop parallelism" } do x = 1, 10 - !$acc loop ! { dg-message "note: Detected parallelism <acc loop vector>" } + !$acc loop ! { dg-message "note: assigned OpenACC worker loop parallelism" } do y = 1, 10 + !$acc loop ! { dg-message "note: assigned OpenACC vector loop parallelism" } + do z = 1, 10 + end do end do end do - !$acc parallel loop gang ! { dg-message "note: Detected parallelism <acc loop gang>" } + !$acc parallel do x = 1, 10 - !$acc loop worker ! { dg-message "note: Detected parallelism <acc loop worker>" } + !$acc loop ! { dg-message "note: assigned OpenACC gang worker loop parallelism" } do y = 1, 10 - !$acc loop vector ! { dg-message "note: Detected parallelism <acc loop vector>" } + !$acc loop ! { dg-message "note: assigned OpenACC vector loop parallelism" } do z = 1, 10 end do end do end do + !$acc end parallel + + !$acc parallel loop seq ! { dg-message "note: assigned OpenACC seq loop parallelism" } + do x = 1, 10 + !$acc loop ! { dg-message "note: assigned OpenACC gang worker loop parallelism" } + do y = 1, 10 + !$acc loop ! { dg-message "note: assigned OpenACC vector loop parallelism" } + do z = 1, 10 + end do + end do + end do + + !$acc parallel loop ! { dg-message "note: assigned OpenACC gang worker loop parallelism" } + do x = 1, 10 + !$acc loop seq ! { dg-message "note: assigned OpenACC seq loop parallelism" } + do y = 1, 10 + !$acc loop ! { dg-message "note: assigned OpenACC vector loop parallelism" } + do z = 1, 10 + end do + end do + end do + + !$acc parallel loop ! { dg-message "note: assigned OpenACC gang worker loop parallelism" } + do x = 1, 10 + !$acc loop ! { dg-message "note: assigned OpenACC vector loop parallelism" } + do y = 1, 10 + !$acc loop seq ! { dg-message "note: assigned OpenACC seq loop parallelism" } + do z = 1, 10 + end do + end do + end do + + !$acc parallel loop seq ! { dg-message "note: assigned OpenACC seq loop parallelism" } + do x = 1, 10 + !$acc loop ! { dg-message "note: assigned OpenACC gang vector loop parallelism" } + do y = 1, 10 + !$acc loop seq ! { dg-message "note: assigned OpenACC seq loop parallelism" } + do z = 1, 10 + end do + end do + end do + end program test diff --git a/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90 b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90 index c10a5838fb0..02c5c2138b5 100644 --- a/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90 @@ -1,14 +1,14 @@ ! Ensure that the middle end does not assign gang level parallelism to ! orphan loop containing reductions. -! { dg-additional-options "-fopt-info-note-omp" } +! { dg-additional-options "-fopt-info-optimized-omp" } subroutine s1 ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" } implicit none !$acc routine gang integer i, sum - !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism <acc loop worker vector>" } + !$acc loop reduction (+:sum) ! { dg-message "note: assigned OpenACC worker vector loop parallelism" } do i = 1, 10 sum = sum + 1 end do @@ -19,9 +19,9 @@ subroutine s2 ! { dg-warning "region is gang partitioned but does not contain ga !$acc routine gang integer i, j, sum - !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism <acc loop worker>" } + !$acc loop reduction (+:sum) ! { dg-message "note: assigned OpenACC worker loop parallelism" } do i = 1, 10 - !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism <acc loop vector>" } + !$acc loop reduction (+:sum) ! { dg-message "note: assigned OpenACC vector loop parallelism" } do j = 1, 10 sum = sum + 1 end do @@ -33,12 +33,12 @@ subroutine s3 ! { dg-warning "region is gang partitioned but does not contain ga !$acc routine gang integer i, j, k, sum - !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism <acc loop worker>" } + !$acc loop reduction (+:sum) ! { dg-message "note: assigned OpenACC worker loop parallelism" } do i = 1, 10 - !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism <acc loop seq>" } + !$acc loop reduction (+:sum) ! { dg-message "note: assigned OpenACC seq loop parallelism" } ! { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } .-1 } do j = 1, 10 - !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism <acc loop vector>" } + !$acc loop reduction (+:sum) ! { dg-message "note: assigned OpenACC vector loop parallelism" } do k = 1, 10 sum = sum + 1 end do @@ -52,16 +52,16 @@ subroutine s4 integer i, j, k, sum !$acc parallel copy(sum) - !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism <acc loop gang vector>" } + !$acc loop reduction (+:sum) ! { dg-message "note: assigned OpenACC gang vector loop parallelism" } do i = 1, 10 sum = sum + 1 end do !$acc end parallel !$acc parallel copy(sum) - !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism <acc loop gang worker>" } + !$acc loop reduction (+:sum) ! { dg-message "note: assigned OpenACC gang worker loop parallelism" } do i = 1, 10 - !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism <acc loop vector>" } + !$acc loop reduction (+:sum) ! { dg-message "note: assigned OpenACC vector loop parallelism" } do j = 1, 10 sum = sum + 1 end do @@ -69,11 +69,11 @@ subroutine s4 !$acc end parallel !$acc parallel copy(sum) - !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism <acc loop gang>" } + !$acc loop reduction (+:sum) ! { dg-message "note: assigned OpenACC gang loop parallelism" } do i = 1, 10 - !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism <acc loop worker>" } + !$acc loop reduction (+:sum) ! { dg-message "note: assigned OpenACC worker loop parallelism" } do j = 1, 10 - !$acc loop reduction (+:sum) ! { dg-message "Detected parallelism <acc loop vector>" } + !$acc loop reduction (+:sum) ! { dg-message "note: assigned OpenACC vector loop parallelism" } do k = 1, 10 sum = sum + 1 end do -- 2.17.1 ^ permalink raw reply [flat|nested] 12+ messages in thread
* [PATCH 2/3] Correct the reported line number in c++ combined OpenACC directives 2018-07-25 15:29 [PATCH 0/3] Add OpenACC diagnostics to -fopt-info-note-omp Cesar Philippidis 2018-07-25 15:29 ` [PATCH 1/3] Correct the reported line number in fortran combined OpenACC directives Cesar Philippidis 2018-07-25 15:29 ` [PATCH 3/3] Add user-friendly OpenACC diagnostics regarding detected parallelism Cesar Philippidis @ 2018-07-25 15:29 ` Cesar Philippidis 2018-12-09 13:02 ` Thomas Schwinge 2 siblings, 1 reply; 12+ messages in thread From: Cesar Philippidis @ 2018-07-25 15:29 UTC (permalink / raw) To: gcc-patches, jakub Like the fortran FE, the C++ FE doesn't set the expr_location of the split acc loop in combined acc parallel/kernels loop directives. This only happens for with combined directives, otherwise cp_parser_omp_construct would be responsible for setting the location. After fixing this bug, I was able to resolve a couple of long standing diagnostics discrepancies between the c/c++ FEs in the test suite. Is this patch OK for trunk? I bootstrapped and regtested using x86_64 with nvptx offloading. Thanks, Cesar 2018-XX-YY Cesar Philippidis <cesar@codesourcery.com> gcc/cp/ * parser.c (cp_parser_oacc_kernels_parallel): Adjust EXPR_LOCATION on the combined acc loop. gcc/testsuite/ * c-c++-common/goacc/combined-directives-3.c: New test. * c-c++-common/goacc/loop-2-kernels.c (void K): Adjust test. * c-c++-common/goacc/loop-2-parallel.c (void P): Adjust test. * c-c++-common/goacc/loop-3.c (void p2): Adjust test. (cherry picked from gomp-4_0-branch r245673) diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 90d5d00..52e61fc 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -37183,8 +37183,9 @@ cp_parser_oacc_kernels_parallel (cp_parser *parser, cp_token *pragma_tok, cp_lexer_consume_token (parser->lexer); tree block = begin_omp_parallel (); tree clauses; - cp_parser_oacc_loop (parser, pragma_tok, p_name, mask, &clauses, - if_p); + tree stmt = cp_parser_oacc_loop (parser, pragma_tok, p_name, mask, + &clauses, if_p); + protected_set_expr_location (stmt, pragma_tok->location); return finish_omp_construct (code, block, clauses); } } diff --git a/gcc/testsuite/c-c++-common/goacc/combined-directives-3.c b/gcc/testsuite/c-c++-common/goacc/combined-directives-3.c new file mode 100644 index 0000000..77d4182 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/combined-directives-3.c @@ -0,0 +1,24 @@ +/* Verify the accuracy of the line number associated with combined + constructs. */ + +int +main () +{ + int x, y, z; + +#pragma acc parallel loop seq auto /* { dg-error "'seq' overrides other OpenACC loop specifiers" } */ + for (x = 0; x < 10; x++) +#pragma acc loop + for (y = 0; y < 10; y++) + ; + +#pragma acc parallel loop gang auto /* { dg-error "'auto' conflicts with other OpenACC loop specifiers" } */ + for (x = 0; x < 10; x++) +#pragma acc loop worker auto /* { dg-error "'auto' conflicts with other OpenACC loop specifiers" } */ + for (y = 0; y < 10; y++) +#pragma acc loop vector + for (z = 0; z < 10; z++) + ; + + return 0; +} 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 01ad32d..3a11ef5f 100644 --- a/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c +++ b/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c @@ -145,8 +145,8 @@ 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" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } } +#pragma acc kernels loop seq worker // { dg-error "'seq' overrides" } + for (i = 0; i < 10; i++) { } #pragma acc kernels loop gang worker for (i = 0; i < 10; i++) @@ -161,8 +161,8 @@ 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" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } } +#pragma acc kernels loop seq vector // { dg-error "'seq' overrides" } + for (i = 0; i < 10; i++) { } #pragma acc kernels loop gang vector for (i = 0; i < 10; i++) @@ -174,16 +174,16 @@ void K(void) #pragma acc kernels loop auto for (i = 0; i < 10; i++) { } -#pragma acc kernels loop seq auto // { dg-error "'seq' overrides" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } } +#pragma acc kernels loop seq auto // { dg-error "'seq' overrides" } + for (i = 0; i < 10; i++) { } -#pragma acc kernels loop gang auto // { dg-error "'auto' conflicts" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } } +#pragma acc kernels loop gang auto // { dg-error "'auto' conflicts" } + for (i = 0; i < 10; i++) { } -#pragma acc kernels loop worker auto // { dg-error "'auto' conflicts" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } } +#pragma acc kernels loop worker auto // { dg-error "'auto' conflicts" } + for (i = 0; i < 10; i++) { } -#pragma acc kernels loop vector auto // { dg-error "'auto' conflicts" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } } +#pragma acc kernels loop vector auto // { dg-error "'auto' conflicts" } + for (i = 0; i < 10; i++) { } } diff --git a/gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c b/gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c index 0ef5741..27f7bbd 100644 --- a/gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c +++ b/gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c @@ -115,16 +115,16 @@ void P(void) for (i = 0; i < 10; i++) { } -#pragma acc parallel loop seq gang // { dg-error "'seq' overrides" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } } +#pragma acc parallel loop seq gang // { dg-error "'seq' overrides" } + for (i = 0; i < 10; i++) { } #pragma acc parallel loop worker for (i = 0; i < 10; i++) { } -#pragma acc parallel loop seq worker // { dg-error "'seq' overrides" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } } +#pragma acc parallel loop seq worker // { dg-error "'seq' overrides" } + for (i = 0; i < 10; i++) { } #pragma acc parallel loop gang worker for (i = 0; i < 10; i++) @@ -134,8 +134,8 @@ void P(void) for (i = 0; i < 10; i++) { } -#pragma acc parallel loop seq vector // { dg-error "'seq' overrides" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } } +#pragma acc parallel loop seq vector // { dg-error "'seq' overrides" } + for (i = 0; i < 10; i++) { } #pragma acc parallel loop gang vector for (i = 0; i < 10; i++) @@ -147,16 +147,16 @@ void P(void) #pragma acc parallel loop auto for (i = 0; i < 10; i++) { } -#pragma acc parallel loop seq auto // { dg-error "'seq' overrides" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } } +#pragma acc parallel loop seq auto // { dg-error "'seq' overrides" } + for (i = 0; i < 10; i++) { } -#pragma acc parallel loop gang auto // { dg-error "'auto' conflicts" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } } +#pragma acc parallel loop gang auto // { dg-error "'auto' conflicts" } + for (i = 0; i < 10; i++) { } -#pragma acc parallel loop worker auto // { dg-error "'auto' conflicts" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } } +#pragma acc parallel loop worker auto // { dg-error "'auto' conflicts" } + for (i = 0; i < 10; i++) { } -#pragma acc parallel loop vector auto // { dg-error "'auto' conflicts" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } } +#pragma acc parallel loop vector auto // { dg-error "'auto' conflicts" } + for (i = 0; i < 10; i++) { } } diff --git a/gcc/testsuite/c-c++-common/goacc/loop-3.c b/gcc/testsuite/c-c++-common/goacc/loop-3.c index 44b65a8..ad31d05 100644 --- a/gcc/testsuite/c-c++-common/goacc/loop-3.c +++ b/gcc/testsuite/c-c++-common/goacc/loop-3.c @@ -35,24 +35,24 @@ void p2 (void) { int i, j; -#pragma acc parallel loop gang(5) // { dg-error "argument not permitted" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } } +#pragma acc parallel loop gang(5) // { dg-error "argument not permitted" } + for (i = 0; i < 10; i++) { } -#pragma acc parallel loop gang(num:5) // { dg-error "argument not permitted" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } } +#pragma acc parallel loop gang(num:5) // { dg-error "argument not permitted" } + for (i = 0; i < 10; i++) { } -#pragma acc parallel loop worker(5) // { dg-error "argument not permitted" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } } +#pragma acc parallel loop worker(5) // { dg-error "argument not permitted" } + for (i = 0; i < 10; i++) { } -#pragma acc parallel loop worker(num:5) // { dg-error "argument not permitted" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } } +#pragma acc parallel loop worker(num:5) // { dg-error "argument not permitted" } + for (i = 0; i < 10; i++) { } -#pragma acc parallel loop vector(5) // { dg-error "argument not permitted" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } } +#pragma acc parallel loop vector(5) // { dg-error "argument not permitted" } + for (i = 0; i < 10; i++) { } -#pragma acc parallel loop vector(length:5) // { dg-error "argument not permitted" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } } +#pragma acc parallel loop vector(length:5) // { dg-error "argument not permitted" } + for (i = 0; i < 10; i++) { } } -- 2.7.4 ^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [PATCH 2/3] Correct the reported line number in c++ combined OpenACC directives 2018-07-25 15:29 ` [PATCH 2/3] Correct the reported line number in c++ combined OpenACC directives Cesar Philippidis @ 2018-12-09 13:02 ` Thomas Schwinge 0 siblings, 0 replies; 12+ messages in thread From: Thomas Schwinge @ 2018-12-09 13:02 UTC (permalink / raw) To: gcc-patches; +Cc: jakub Hi! On Wed, 25 Jul 2018 08:29:18 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote: > Like the fortran FE, the C++ FE doesn't set the expr_location of the > split acc loop in combined acc parallel/kernels loop directives. This > only happens for with combined directives, otherwise > cp_parser_omp_construct would be responsible for setting the > location. After fixing this bug, I was able to resolve a couple of > long standing diagnostics discrepancies between the c/c++ FEs in the > test suite. > > Is this patch OK for trunk? I bootstrapped and regtested using x86_64 > with nvptx offloading. Thanks, committed to trunk in r266923: commit 21fb940c5d500a1a8e850b59412fb506aa51181a Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> Date: Sun Dec 9 12:48:26 2018 +0000 Correct the reported line number in C++ combined OpenACC directives The C++ FE doesn't set the expr_location of the split acc loop in combined acc parallel/kernels loop directives. This only happens for with combined directives, otherwise cp_parser_omp_construct would be responsible for setting the location. After fixing this bug, I was able to resolve a couple of long standing diagnostics discrepancies between the C/C++ FEs in the test suite. gcc/cp/ * parser.c (cp_parser_oacc_kernels_parallel): Adjust EXPR_LOCATION on the combined acc loop. gcc/testsuite/ * c-c++-common/goacc/combined-directives-3.c: New test. * c-c++-common/goacc/loop-2-kernels.c (void K): Adjust test. * c-c++-common/goacc/loop-2-parallel.c (void P): Adjust test. * c-c++-common/goacc/loop-3.c (void p2): Adjust test. Reviewed-by: Thomas Schwinge <thomas@codesourcery.com> git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@266923 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/cp/ChangeLog | 5 ++++ gcc/cp/parser.c | 5 ++-- gcc/testsuite/ChangeLog | 7 ++++++ .../c-c++-common/goacc/combined-directives-3.c | 24 +++++++++++++++++++ gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c | 24 +++++++++---------- gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c | 28 +++++++++++----------- gcc/testsuite/c-c++-common/goacc/loop-3.c | 24 +++++++++---------- 7 files changed, 77 insertions(+), 40 deletions(-) diff --git gcc/cp/ChangeLog gcc/cp/ChangeLog index fc22f206606b..672be2d072db 100644 --- gcc/cp/ChangeLog +++ gcc/cp/ChangeLog @@ -1,3 +1,8 @@ +2018-12-09 Cesar Philippidis <cesar@codesourcery.com> + + * parser.c (cp_parser_oacc_kernels_parallel): Adjust EXPR_LOCATION + on the combined acc loop. + 2018-12-07 Paolo Carlini <paolo.carlini@oracle.com> * decl2.c (grokbitfield): Use DECL_SOURCE_LOCATION in error messages diff --git gcc/cp/parser.c gcc/cp/parser.c index adfe09e494dc..8b669a82b147 100644 --- gcc/cp/parser.c +++ gcc/cp/parser.c @@ -38742,8 +38742,9 @@ cp_parser_oacc_kernels_parallel (cp_parser *parser, cp_token *pragma_tok, cp_lexer_consume_token (parser->lexer); tree block = begin_omp_parallel (); tree clauses; - cp_parser_oacc_loop (parser, pragma_tok, p_name, mask, &clauses, - if_p); + tree stmt = cp_parser_oacc_loop (parser, pragma_tok, p_name, mask, + &clauses, if_p); + protected_set_expr_location (stmt, pragma_tok->location); return finish_omp_construct (code, block, clauses); } } diff --git gcc/testsuite/ChangeLog gcc/testsuite/ChangeLog index 192a29ee971c..6b26f6f510db 100644 --- gcc/testsuite/ChangeLog +++ gcc/testsuite/ChangeLog @@ -1,3 +1,10 @@ +2018-12-09 Cesar Philippidis <cesar@codesourcery.com> + + * c-c++-common/goacc/combined-directives-3.c: New test. + * c-c++-common/goacc/loop-2-kernels.c (void K): Adjust test. + * c-c++-common/goacc/loop-2-parallel.c (void P): Adjust test. + * c-c++-common/goacc/loop-3.c (void p2): Adjust test. + 2018-12-09 Thomas Schwinge <thomas@codesourcery.com> * gfortran.dg/goacc/loop-2-kernels.f95: Update. diff --git gcc/testsuite/c-c++-common/goacc/combined-directives-3.c gcc/testsuite/c-c++-common/goacc/combined-directives-3.c new file mode 100644 index 000000000000..77d418262eac --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/combined-directives-3.c @@ -0,0 +1,24 @@ +/* Verify the accuracy of the line number associated with combined + constructs. */ + +int +main () +{ + int x, y, z; + +#pragma acc parallel loop seq auto /* { dg-error "'seq' overrides other OpenACC loop specifiers" } */ + for (x = 0; x < 10; x++) +#pragma acc loop + for (y = 0; y < 10; y++) + ; + +#pragma acc parallel loop gang auto /* { dg-error "'auto' conflicts with other OpenACC loop specifiers" } */ + for (x = 0; x < 10; x++) +#pragma acc loop worker auto /* { dg-error "'auto' conflicts with other OpenACC loop specifiers" } */ + for (y = 0; y < 10; y++) +#pragma acc loop vector + for (z = 0; z < 10; z++) + ; + + return 0; +} diff --git gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c index 93e1cece26b4..01515089a7d8 100644 --- gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c +++ gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c @@ -147,8 +147,8 @@ 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" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } } +#pragma acc kernels loop seq worker // { dg-error "'seq' overrides" } + for (i = 0; i < 10; i++) { } #pragma acc kernels loop gang worker for (i = 0; i < 10; i++) @@ -163,8 +163,8 @@ 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" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } } +#pragma acc kernels loop seq vector // { dg-error "'seq' overrides" } + for (i = 0; i < 10; i++) { } #pragma acc kernels loop gang vector for (i = 0; i < 10; i++) @@ -176,16 +176,16 @@ void K(void) #pragma acc kernels loop auto for (i = 0; i < 10; i++) { } -#pragma acc kernels loop seq auto // { dg-error "'seq' overrides" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } } +#pragma acc kernels loop seq auto // { dg-error "'seq' overrides" } + for (i = 0; i < 10; i++) { } -#pragma acc kernels loop gang auto // { dg-error "'auto' conflicts" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } } +#pragma acc kernels loop gang auto // { dg-error "'auto' conflicts" } + for (i = 0; i < 10; i++) { } -#pragma acc kernels loop worker auto // { dg-error "'auto' conflicts" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } } +#pragma acc kernels loop worker auto // { dg-error "'auto' conflicts" } + for (i = 0; i < 10; i++) { } -#pragma acc kernels loop vector auto // { dg-error "'auto' conflicts" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } } +#pragma acc kernels loop vector auto // { dg-error "'auto' conflicts" } + for (i = 0; i < 10; i++) { } } diff --git gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c index 5b1e9d7ce87f..e3e6786e572c 100644 --- gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c +++ gcc/testsuite/c-c++-common/goacc/loop-2-parallel.c @@ -117,16 +117,16 @@ void P(void) for (i = 0; i < 10; i++) { } -#pragma acc parallel loop seq gang // { dg-error "'seq' overrides" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } } +#pragma acc parallel loop seq gang // { dg-error "'seq' overrides" } + for (i = 0; i < 10; i++) { } #pragma acc parallel loop worker for (i = 0; i < 10; i++) { } -#pragma acc parallel loop seq worker // { dg-error "'seq' overrides" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } } +#pragma acc parallel loop seq worker // { dg-error "'seq' overrides" } + for (i = 0; i < 10; i++) { } #pragma acc parallel loop gang worker for (i = 0; i < 10; i++) @@ -136,8 +136,8 @@ void P(void) for (i = 0; i < 10; i++) { } -#pragma acc parallel loop seq vector // { dg-error "'seq' overrides" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } } +#pragma acc parallel loop seq vector // { dg-error "'seq' overrides" } + for (i = 0; i < 10; i++) { } #pragma acc parallel loop gang vector for (i = 0; i < 10; i++) @@ -149,16 +149,16 @@ void P(void) #pragma acc parallel loop auto for (i = 0; i < 10; i++) { } -#pragma acc parallel loop seq auto // { dg-error "'seq' overrides" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } } +#pragma acc parallel loop seq auto // { dg-error "'seq' overrides" } + for (i = 0; i < 10; i++) { } -#pragma acc parallel loop gang auto // { dg-error "'auto' conflicts" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } } +#pragma acc parallel loop gang auto // { dg-error "'auto' conflicts" } + for (i = 0; i < 10; i++) { } -#pragma acc parallel loop worker auto // { dg-error "'auto' conflicts" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } } +#pragma acc parallel loop worker auto // { dg-error "'auto' conflicts" } + for (i = 0; i < 10; i++) { } -#pragma acc parallel loop vector auto // { dg-error "'auto' conflicts" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } } +#pragma acc parallel loop vector auto // { dg-error "'auto' conflicts" } + for (i = 0; i < 10; i++) { } } diff --git gcc/testsuite/c-c++-common/goacc/loop-3.c gcc/testsuite/c-c++-common/goacc/loop-3.c index e6c3f18042e6..ad5a4bd5aa9a 100644 --- gcc/testsuite/c-c++-common/goacc/loop-3.c +++ gcc/testsuite/c-c++-common/goacc/loop-3.c @@ -37,24 +37,24 @@ void p2 (void) { int i, j; -#pragma acc parallel loop gang(5) // { dg-error "argument not permitted" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } } +#pragma acc parallel loop gang(5) // { dg-error "argument not permitted" } + for (i = 0; i < 10; i++) { } -#pragma acc parallel loop gang(num:5) // { dg-error "argument not permitted" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } } +#pragma acc parallel loop gang(num:5) // { dg-error "argument not permitted" } + for (i = 0; i < 10; i++) { } -#pragma acc parallel loop worker(5) // { dg-error "argument not permitted" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } } +#pragma acc parallel loop worker(5) // { dg-error "argument not permitted" } + for (i = 0; i < 10; i++) { } -#pragma acc parallel loop worker(num:5) // { dg-error "argument not permitted" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } } +#pragma acc parallel loop worker(num:5) // { dg-error "argument not permitted" } + for (i = 0; i < 10; i++) { } -#pragma acc parallel loop vector(5) // { dg-error "argument not permitted" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } } +#pragma acc parallel loop vector(5) // { dg-error "argument not permitted" } + for (i = 0; i < 10; i++) { } -#pragma acc parallel loop vector(length:5) // { dg-error "argument not permitted" "" { target c } } - for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } } +#pragma acc parallel loop vector(length:5) // { dg-error "argument not permitted" } + for (i = 0; i < 10; i++) { } } Grüße Thomas ^ permalink raw reply [flat|nested] 12+ messages in thread
end of thread, other threads:[~2019-01-28 21:35 UTC | newest] Thread overview: 12+ messages (download: mbox.gz / follow: Atom feed) -- links below jump to the message on this page -- 2018-07-25 15:29 [PATCH 0/3] Add OpenACC diagnostics to -fopt-info-note-omp Cesar Philippidis 2018-07-25 15:29 ` [PATCH 1/3] Correct the reported line number in fortran combined OpenACC directives Cesar Philippidis 2018-07-25 15:32 ` Marek Polacek 2018-07-25 15:53 ` Cesar Philippidis 2018-12-09 13:07 ` Thomas Schwinge 2018-07-25 15:29 ` [PATCH 3/3] Add user-friendly OpenACC diagnostics regarding detected parallelism Cesar Philippidis 2018-07-26 8:34 ` Richard Biener 2018-07-26 14:14 ` Cesar Philippidis 2018-12-14 21:03 ` Add user-friendly diagnostics for OpenACC loop parallelism assigned (was: [PATCH 3/3] Add user-friendly OpenACC diagnostics regarding detected parallelism) Thomas Schwinge 2019-01-28 21:51 ` [og8] Add user-friendly diagnostics for OpenACC loop parallelism assigned Thomas Schwinge 2018-07-25 15:29 ` [PATCH 2/3] Correct the reported line number in c++ combined OpenACC directives Cesar Philippidis 2018-12-09 13:02 ` Thomas Schwinge
This is a public inbox, see mirroring instructions for how to clone and mirror all data and code used for this inbox; as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).