This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

Re: [PATCH 3/3] Add user-friendly OpenACC diagnostics regarding detected parallelism.


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
>


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]