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]

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


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


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