Annotate inner loops in "acc kernels loop" directives (Fortran).

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 in Fortran.

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

	gcc/fortran/
	* openmp.cc (annotate_do_loops_in_kernels): Handle
	EXEC_OACC_KERNELS_LOOP separately to permit annotation of inner
	loops in a combined "acc kernels loop" directive.

	gcc/testsuite/
	* gfortran.dg/goacc/kernels-loop-annotation-18.f95: New.
	* gfortran.dg/goacc/kernels-loop-annotation-19.f95: New.
	* gfortran.dg/goacc/combined-directives.f90: Adjust expected
	patterns.
	* gfortran.dg/goacc/private-explicit-kernels-1.f95: Likewise.
	* gfortran.dg/goacc/private-predetermined-kernels-1.f95:
	Likewise.

diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index df352c58706e55a445761b5b04637c01a2158969..e2a8b51f7ecb35aaf55cfa3ebf702a61f4c208e5 100644 (file)
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,11 @@
+2020-08-19  Sandra Loosemore  <sandra@codesourcery.com>
+
+       Annotate inner loops in "acc kernels loop" directives (Fortran).
+
+       * openmp.cc (annotate_do_loops_in_kernels): Handle
+       EXEC_OACC_KERNELS_LOOP separately to permit annotation of inner
+       loops in a combined "acc kernels loop" directive.
+
 2020-08-19  Sandra Loosemore  <sandra@codesourcery.com>
 
        Add a "combined" flag for "acc kernels loop" etc directives.

diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 387ef866557a60a1e9e64a3a4c0af89539dab9da..09804e71b82c77faff6a148490f6b012f89c2227 100644 (file)
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -10562,7 +10562,6 @@ annotate_do_loops_in_kernels (gfc_code *code, gfc_code *parent,
 
        case EXEC_OACC_PARALLEL_LOOP:
        case EXEC_OACC_PARALLEL:
-       case EXEC_OACC_KERNELS_LOOP:
        case EXEC_OACC_LOOP:
          /* Do not try to add automatic OpenACC annotations inside manually
             annotated loops.  Presumably, the user avoided doing it on
@@ -10607,6 +10606,55 @@ annotate_do_loops_in_kernels (gfc_code *code, gfc_code *parent,
            }
          break;
 
+       case EXEC_OACC_KERNELS_LOOP:
+         /* This is a combined "acc kernels loop" directive.  We want to
+            leave the outer loop alone but try to annotate any nested
+            loops in the body.  The expected structure nesting here is
+              EXEC_OACC_KERNELS_LOOP
+                EXEC_OACC_KERNELS_LOOP
+                  EXEC_DO
+                    EXEC_DO
+                      ...body...  */
+         if (code->block)
+           /* Might be empty?  */
+           {
+             gcc_assert (code->block->op == EXEC_OACC_KERNELS_LOOP);
+             gfc_omp_clauses *clauses = code->ext.omp_clauses;
+             int collapse = clauses->collapse;
+             gfc_expr_list *tile = clauses->tile_list;
+             gfc_code *inner = code->block->next;
+
+             gcc_assert (inner->op == EXEC_DO);
+             gcc_assert (inner->block->op == EXEC_DO);
+
+             /* We need to skip over nested loops covered by "collapse" or
+                "tile" clauses.  "Tile" takes precedence
+                (see gfc_trans_omp_do).  */
+             if (tile)
+               {
+                 collapse = 0;
+                 for (gfc_expr_list *el = tile; el; el = el->next)
+                   collapse++;
+               }
+             if (clauses->orderedc)
+               collapse = clauses->orderedc;
+             if (collapse <= 0)
+               collapse = 1;
+             for (int i = 1; i < collapse; i++)
+               {
+                 gcc_assert (inner->op == EXEC_DO);
+                 gcc_assert (inner->block->op == EXEC_DO);
+                 inner = inner->block->next;
+               }
+             if (inner)
+               /* Loop might have empty body?  */
+               annotate_do_loops_in_kernels (inner->block->next,
+                                             inner, goto_targets,
+                                             as_in_kernels_region);
+           }
+         walk_block = false;
+         break;
+
        case EXEC_DO_WHILE:
        case EXEC_DO_CONCURRENT:
          /* Traverse the body in a special state to allow EXIT statements

diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index ff201756e51576fe10fb46b9f2d0ea32879a7ca5..8f4aa70f07411603ed9092fa9a4b8908767d60b5 100644 (file)
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,5 +1,17 @@
 2020-08-19   Sandra Loosemore  <sandra@codesourcery.com>
 
+       Annotate inner loops in "acc kernels loop" directives (Fortran).
+
+       * gfortran.dg/goacc/kernels-loop-annotation-18.f95: New.
+       * gfortran.dg/goacc/kernels-loop-annotation-19.f95: New.
+       * gfortran.dg/goacc/combined-directives.f90: Adjust expected
+       patterns.
+       * gfortran.dg/goacc/private-explicit-kernels-1.f95: Likewise.
+       * gfortran.dg/goacc/private-predetermined-kernels-1.f95:
+       Likewise.
+
+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.

diff --git a/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
index 956349204f4d2353648adbe71a407400b92e6e5f..562a4e40cd7d9141a070677f0850b794a15252f9 100644 (file)
--- a/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
@@ -139,10 +139,21 @@ end subroutine test
 
 ! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. collapse.2." 2 "gimple" } }
 ! { dg-final { scan-tree-dump-times "acc loop private.i. gang" 2 "gimple" } }
-! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. worker" 2 "gimple" } }
-! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. vector" 2 "gimple" } }
-! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. seq" 2 "gimple" } }
-! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. auto" 2 "gimple" } }
+
+! These are the parallel loop variants.
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. worker" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. vector" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. seq" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. auto" 1 "gimple" } }
+
+! These are the kernels loop variants.  Here the inner loops are annotated
+! separately.
+! { dg-final { scan-tree-dump-times "acc loop private.i. worker" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. vector" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. seq" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop private.i. auto" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "acc loop auto private.j." 4 "gimple" } }
+
 ! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. tile.2, 3" 2 "gimple" } }
 ! { dg-final { scan-tree-dump-times "acc loop private.i. independent" 2 "gimple" } }
 ! { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } }

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-18.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-18.f95
new file mode 100644 (file)
index 0000000..e4e210a
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-18.f95
@@ -0,0 +1,28 @@
+! { 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.
+
+subroutine f (a, b)
+
+  implicit none
+  real, intent (in), dimension(20) :: a
+  real, intent (out), dimension(20) :: b
+  integer :: k, l, m
+
+!$acc kernels loop
+  do k = 1, 20
+    do l = 1, 20
+      do m = 1, 20
+       b(m) = a(m);
+      end do
+    end do
+  end do
+
+end subroutine f
+
+! { dg-final { scan-tree-dump-times "acc loop auto" 2 "original" } }
+

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-19.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-19.f95
new file mode 100644 (file)
index 0000000..5dd6e7f
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-19.f95
@@ -0,0 +1,29 @@
+! { 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.
+
+subroutine f (a, b)
+
+  implicit none
+  real, intent (in), dimension(20) :: a
+  real, intent (out), dimension(20) :: b
+  integer :: k, l, m
+
+!$acc kernels loop collapse(2)
+  do k = 1, 20
+    do l = 1, 20
+      do m = 1, 20
+       b(m) = a(m);
+      end do
+    end do
+  end do
+
+end subroutine f
+
+! { dg-final { scan-tree-dump-times "acc loop .*collapse.2." 1 "original" } }
+! { dg-final { scan-tree-dump-times "acc loop auto" 1 "original" } }
+

diff --git a/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
index 5d563d226b0c30439c62713e113c4a3ec139fda3..0c47045df9c82f507bd3287a0223fb5e47f7441c 100644 (file)
--- a/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
@@ -73,8 +73,9 @@ program test
 
   !$acc kernels loop private(i2_1_c, j2_1_c) independent
   ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\) private\\(j2_1_c\\) independent" 1 "original" } }
-  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\) private\\(j2_1_c\\) independent" 1 "gimple" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\) independent" 1 "gimple" } }
   do i2_1_c = 1, 100
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop auto private\\(j2_1_c\\)" 1 "gimple" } }
      do j2_1_c = 1, 100
      end do
   end do
@@ -130,9 +131,11 @@ program test
 
   !$acc kernels loop private(i3_1_c, j3_1_c, k3_1_c) independent
   ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\) private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "original" } }
-  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\) private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "gimple" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\) independent" 1 "gimple" } }
   do i3_1_c = 1, 100
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop auto private\\(j3_1_c\\)" 1 "gimple" } }
      do j3_1_c = 1, 100
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop auto private\\(k3_1_c\\)" 1 "gimple" } }
         do k3_1_c = 1, 100
         end do
      end do

diff --git a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
index 12a7854526a94a950ccc9b1facdb11bc20d9c451..3357a20263e702fb6d599d773d38f3876e9d9fa5 100644 (file)
--- a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
@@ -73,8 +73,9 @@ program test
 
   !$acc kernels loop independent
   ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\) private\\(j2_1_c\\) independent" 1 "original" } }
-  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\) private\\(j2_1_c\\) independent" 1 "gimple" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i2_1_c\\) independent" 1 "gimple" } }
   do i2_1_c = 1, 100
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop auto private\\(j2_1_c\\)" 1 "gimple" } }
      do j2_1_c = 1, 100
      end do
   end do
@@ -130,9 +131,11 @@ program test
 
   !$acc kernels loop independent
   ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\) private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "original" } }
-  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\) private\\(j3_1_c\\) private\\(k3_1_c\\) independent" 1 "gimple" } }
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(i3_1_c\\) independent" 1 "gimple" } }
   do i3_1_c = 1, 100
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop auto private\\(j3_1_c\\)" 1 "gimple" } }
      do j3_1_c = 1, 100
+  ! { dg-final { scan-tree-dump-times "#pragma acc loop auto private\\(k3_1_c\\)" 1 "gimple" } }
         do k3_1_c = 1, 100
         end do
      end do