aboutsummaryrefslogtreecommitdiff
path: root/gcc
diff options
context:
space:
mode:
Diffstat (limited to 'gcc')
-rw-r--r--gcc/fortran/ChangeLog.omp8
-rw-r--r--gcc/fortran/openmp.cc50
-rw-r--r--gcc/testsuite/ChangeLog.omp12
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/combined-directives.f9019
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-18.f9528
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-19.f9529
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f957
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f957
8 files changed, 151 insertions, 9 deletions
diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index df352c5..e2a8b51 100644
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,5 +1,13 @@
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.
* trans-openmp.cc (gfc_trans_omp_do): Add combined parameter,
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 39a4e65..7b0a593 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -9646,7 +9646,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
@@ -9691,6 +9690,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 54739a1..1e2ed0b 100644
--- 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 9563492..562a4e4 100644
--- 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
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
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 5d563d2..0c47045 100644
--- 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 12a7854..3357a20 100644
--- 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
generated by cgit v1.1 at 2025-09-27 07:43:32 +0000