Annotate inner loops in "acc kernels loop" directives (C/C++).

Normally explicit loop directives in a kernels region inhibit
automatic annotation of other loops in the same nest, on the theory
that users have indicated they want manual control over that section
of code.  However there seems to be an expectation in user code that
the combined "kernels loop" directive should still allow annotation of
inner loops.  This patch implements this behavior for C and C++.

2020-08-19  Sandra Loosemore  <sandra@codesourcery.com>

	gcc/c-family/
	* c-omp.c (annotate_loops_in_kernels_regions): Process inner
	loops in combined "acc kernels loop" directives.

	gcc/testsuite/
	* c-c++-common/goacc/kernels-loop-annotation-18.c: New.
	* c-c++-common/goacc/kernels-loop-annotation-19.c: New.
	* c-c++-common/goacc/combined-directives.c: Adjust expected
	patterns.

diff --git a/gcc/c-family/ChangeLog.omp b/gcc/c-family/ChangeLog.omp
index 92ffa851d4482986adf9e51f7bb632ec382add35..51a3b9e476655c02c05260e3abea1080ca1a271d 100644 (file)
--- a/gcc/c-family/ChangeLog.omp
+++ b/gcc/c-family/ChangeLog.omp
@@ -1,3 +1,10 @@
+2020-08-19  Sandra Loosemore  <sandra@codesourcery.com>
+
+       Annotate inner loops in "acc kernels loop" directives (C/C++).
+
+       * c-omp.c (annotate_loops_in_kernels_regions): Process inner
+       loops in combined "acc kernels loop" directives.
+
 2020-08-18  Kwok Cheung Yeung  <kcy@codesourcery.com>
 
        Backport from mainline

diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index c1d6afa2ca7c18e552a0adf30cb041983cc54e13..24f2448b235e8053a6cba0cd363451c5601de50a 100644 (file)
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -2782,18 +2782,30 @@ annotate_loops_in_kernels_regions (tree *nodeptr, int *walk_subtrees,
       /* Do not try to add automatic OpenACC annotations inside manually
         annotated loops.  Presumably, the user avoided doing it on
         purpose; for example, all available levels of parallelism may
-        have been used up.  */
-      {
-       struct annotation_info nested_info
-         = { NULL_TREE, NULL_TREE, false, as_explicit_annotation,
-             node, info };
-       if (info->state >= as_in_kernels_region)
-         do_not_annotate_loop_nest (info, as_explicit_annotation,
-                                    node);
-       walk_tree (&OMP_BODY (node), annotate_loops_in_kernels_regions,
-                  (void *) &nested_info, NULL);
-       *walk_subtrees = 0;
-      }
+        have been used up.  However, assume that the combined construct
+        "#pragma acc kernels loop" means to try to process the whole
+        loop nest.
+        Note that a single OACC_LOOP construct represents an entire set
+        of collapsed loops so we do not have to deal explicitly with the
+        collapse clause here, as the Fortran front end does.  */
+      if (info->state == as_in_kernels_region && OACC_LOOP_COMBINED (node))
+       {
+         walk_tree (&OMP_BODY (node), annotate_loops_in_kernels_regions,
+                    (void *) info, NULL);
+         *walk_subtrees = 0;
+       }
+      else
+       {
+         struct annotation_info nested_info
+           = { NULL_TREE, NULL_TREE, false, as_explicit_annotation,
+               node, info };
+         if (info->state >= as_in_kernels_region)
+           do_not_annotate_loop_nest (info, as_explicit_annotation,
+                                      node);
+         walk_tree (&OMP_BODY (node), annotate_loops_in_kernels_regions,
+                    (void *) &nested_info, NULL);
+         *walk_subtrees = 0;
+       }
       break;
 
     case FOR_STMT:

diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index ffc8f634f96f9bfe8ca636d5341bf42927e2b427..970345bcad90d09306b8b9d118f82ad40ee9d1aa 100644 (file)
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,12 @@
+2020-08-19   Sandra Loosemore  <sandra@codesourcery.com>
+
+       Annotate inner loops in "acc kernels loop" directives (C/C++).
+
+       * c-c++-common/goacc/kernels-loop-annotation-18.c: New.
+       * c-c++-common/goacc/kernels-loop-annotation-19.c: New.
+       * c-c++-common/goacc/combined-directives.c: Adjust expected
+       patterns.
+
 2020-08-19  Kwok Cheung Yeung  <kcy@codesourcery.com>
 
        * gfortran.dg/goacc/pr70828.f90: Update expected output in Gimple

diff --git a/gcc/testsuite/c-c++-common/goacc/combined-directives.c b/gcc/testsuite/c-c++-common/goacc/combined-directives.c
index c2a3c57b48b83432c1bf000d279a2d09a8f2f57a..2519f23d49f0b3de2faf25acf830d85737ac8754 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/combined-directives.c
+++ b/gcc/testsuite/c-c++-common/goacc/combined-directives.c
@@ -110,7 +110,7 @@ test ()
 // { dg-final { scan-tree-dump-times "acc loop worker" 2 "gimple" } }
 // { dg-final { scan-tree-dump-times "acc loop vector" 2 "gimple" } }
 // { dg-final { scan-tree-dump-times "acc loop seq" 2 "gimple" } }
-// { dg-final { scan-tree-dump-times "acc loop auto" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "acc loop auto" 6 "gimple" } }
 // { dg-final { scan-tree-dump-times "acc loop tile.2, 3" 2 "gimple" } }
 // { dg-final { scan-tree-dump-times "acc loop independent private.i" 2 "gimple" } }
 // { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } }

diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-18.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-18.c
new file mode 100644 (file)
index 0000000..89ec644
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-18.c
@@ -0,0 +1,18 @@
+/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */
+/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */
+/* { dg-additional-options "-fdump-tree-original" } */
+/* { dg-do compile } */
+
+/* Test that "acc kernels loop" directive causes annotation of the entire
+   loop nest.  */
+
+void f (float *a, float *b)
+{
+#pragma acc kernels loop
+  for (int k = 0; k < 20; k++)
+    for (int l = 0; l < 20; l++)
+      for (int m = 0; m < 20; m++)
+       b[m] = a[m];
+}
+
+/* { dg-final { scan-tree-dump-times "acc loop auto" 2 "original" } } */

diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-19.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-19.c
new file mode 100644 (file)
index 0000000..77a3b7a
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-19.c
@@ -0,0 +1,19 @@
+/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */
+/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */
+/* { dg-additional-options "-fdump-tree-original" } */
+/* { dg-do compile } */
+
+/* Test that "acc kernels loop" directive causes annotation of the entire
+   loop nest in the presence of a collapse clause.  */
+
+void f (float *a, float *b)
+{
+#pragma acc kernels loop collapse(2)
+  for (int k = 0; k < 20; k++)
+    for (int l = 0; l < 20; l++)
+      for (int m = 0; m < 20; m++)
+       b[m] = a[m];
+}
+
+/* { dg-final { scan-tree-dump-times "acc loop collapse.2." 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "acc loop auto" 1 "original" } } */