On Wed, Jul 25, 2018 at 5:30 PM Cesar Philippidis
<ce...@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 <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
>
