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 <ce...@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