[committed,3/3,OG10] Annotate inner loops in "acc kernels loop" directives (Fortran).

Message ID	20200820210711.8017-4-sandra@codesourcery.com
State	New
Headers	show Return-Path: <gcc-patches-bounces@gcc.gnu.org> DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org E188B3857C4D IronPort-SDR: 3DA0ZafvEgstTIB2XWlCESyscPbbwFjLwLoB4dUITe5Kg/jbt/oh85o42NqH3pzTL+VLoMzLFf y8urrDAc9r6O1JE+Pf90fnYJPQ8zd9ZtO3F9lRlsrvxRAwE9JyWB8zBktFp4Bpvvb7ubBJ5Zs6 SmckWBdD9uJxpJ+m9e3cAE6Ezrw6X5YRNu/efEYmBu8QPFjB+vIx3XvLUN7gktug4oa/GgaPor tC8R0DTW2aXB4iu5P5kkeM+iyEDjkG/mkJLAN6G+2pSPDcvtDd6aKMy4F4qPHyxyohX2oMgBH4 xNs= IronPort-SDR: fAy53yvTE/BSF+VxEdFcBc/LUX9VUB5KmVSf+wmoPyRhQH7f5eeYqV5S7G4GP4KwaKejN9GjjZ gmEqWTo5ErllioMJzXAz+H++0VF7f8ZtUVJXHzxVVBjAP5Kx+UHO+znFIW4egp5wznDj8EFlg9 SSIcK4few5trozY941ekTAf9VBHyGwUHCXQ5I+Kr0MfCyME6/l0mr/APYQWGSHxrYMESVqxUVj RIIb84eleu9rTxjUwcYNkxfoW+CkmMwDs1D5I7pb7Yzu1MM/xoDLBFmwYTKH33klvH5NCKhLlu XfY= From: Sandra Loosemore <sandra@codesourcery.com> To: <gcc-patches@gcc.gnu.org> Subject: [committed 3/3] [OG10] Annotate inner loops in "acc kernels loop" directives (Fortran). Date: Thu, 20 Aug 2020 15:07:11 -0600 Message-ID: <20200820210711.8017-4-sandra@codesourcery.com> In-Reply-To: <20200820210711.8017-1-sandra@codesourcery.com> References: <20200820210711.8017-1-sandra@codesourcery.com> MIME-Version: 1.0 Content-Type: text/plain Precedence: list Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" <gcc-patches-bounces@gcc.gnu.org>
Series	openacc: Fix annotation of inner loops in combined "acc kernels loop" directives \| expand [committed,0/3,OG10] openacc: Fix annotation of inner loops in combined "acc kernels loop" directiv… [committed,1/3,OG10] Add a "combined" flag for "acc kernels loop" etc directives. [committed,2/3,og10] Annotate inner loops in "acc kernels loop" directives (C/C++). [committed,3/3,OG10] Annotate inner loops in "acc kernels loop" directives (Fortran).

diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp index 1d1ee9e..88b2729 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.c (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.c (gfc_trans_omp_do): Add combined parameter, diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 9d13863..b9e4bda 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -6910,7 +6910,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 @@ -6955,6 +6954,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 970345b..887a24d 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

[committed,3/3,OG10] Annotate inner loops in "acc kernels loop" directives (Fortran).

Commit Message

Patch