[gcc/devel/omp/gcc-12] Annotate inner loops in "acc kernels loop" directives (C/C++).

[gcc/devel/omp/gcc-12] Annotate inner loops in "acc kernels loop" directives (C/C++).

[Kwok Yeung]

https://gcc.gnu.org/g:cdd128b106cd91bf4764e91bdde5b7a8220fd4d0

commit cdd128b106cd91bf4764e91bdde5b7a8220fd4d0
Author: Sandra Loosemore <sandra@codesourcery.com>
Date:   Wed Aug 19 19:18:57 2020 -0700

    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.cc (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:
---
 gcc/c-family/ChangeLog.omp                         |  7 +++++
 gcc/c-family/c-omp.cc                              | 36 ++++++++++++++--------
 gcc/testsuite/ChangeLog.omp                        |  9 ++++++
 .../c-c++-common/goacc/combined-directives.c       |  2 +-
 .../goacc/kernels-loop-annotation-18.c             | 18 +++++++++++
 .../goacc/kernels-loop-annotation-19.c             | 19 ++++++++++++
 6 files changed, 78 insertions(+), 13 deletions(-)

diff --git a/gcc/c-family/ChangeLog.omp b/gcc/c-family/ChangeLog.omp
index 5da3b329dc8..3961227eb2a 100644
--- 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.cc (annotate_loops_in_kernels_regions): Process inner
+	loops in combined "acc kernels loop" directives.
+
 2020-03-27  Sandra Loosemore  <sandra@codesourcery.com>
 
 	* c-common.h (c_oacc_annotate_loops_in_kernels_regions): Declare.
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 7fd40f3e4d9..7d2b69f3632 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -3472,18 +3472,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 a7884d442a0..54739a146bb 100644
--- 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 c2a3c57b48b..2519f23d49f 100644
--- 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
index 00000000000..89ec6447625
--- /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
index 00000000000..77a3b7a9136
--- /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" } } */