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

[Cesar Philippidis <cesar@codesourcery.com>]

[-- 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