From patchwork Thu Aug 20 21:07:11 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Sandra Loosemore X-Patchwork-Id: 1348628 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4BXcgC5rwDz9sR4 for ; Fri, 21 Aug 2020 07:07:51 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 5543F386EC32; Thu, 20 Aug 2020 21:07:42 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id E188B3857C4D for ; Thu, 20 Aug 2020 21:07:38 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org E188B3857C4D Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Sandra_Loosemore@mentor.com IronPort-SDR: 3DA0ZafvEgstTIB2XWlCESyscPbbwFjLwLoB4dUITe5Kg/jbt/oh85o42NqH3pzTL+VLoMzLFf y8urrDAc9r6O1JE+Pf90fnYJPQ8zd9ZtO3F9lRlsrvxRAwE9JyWB8zBktFp4Bpvvb7ubBJ5Zs6 SmckWBdD9uJxpJ+m9e3cAE6Ezrw6X5YRNu/efEYmBu8QPFjB+vIx3XvLUN7gktug4oa/GgaPor tC8R0DTW2aXB4iu5P5kkeM+iyEDjkG/mkJLAN6G+2pSPDcvtDd6aKMy4F4qPHyxyohX2oMgBH4 xNs= X-IronPort-AV: E=Sophos;i="5.76,334,1592899200"; d="scan'208";a="52101887" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 20 Aug 2020 13:07:38 -0800 IronPort-SDR: fAy53yvTE/BSF+VxEdFcBc/LUX9VUB5KmVSf+wmoPyRhQH7f5eeYqV5S7G4GP4KwaKejN9GjjZ gmEqWTo5ErllioMJzXAz+H++0VF7f8ZtUVJXHzxVVBjAP5Kx+UHO+znFIW4egp5wznDj8EFlg9 SSIcK4few5trozY941ekTAf9VBHyGwUHCXQ5I+Kr0MfCyME6/l0mr/APYQWGSHxrYMESVqxUVj RIIb84eleu9rTxjUwcYNkxfoW+CkmMwDs1D5I7pb7Yzu1MM/xoDLBFmwYTKH33klvH5NCKhLlu XfY= From: Sandra Loosemore To: 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> X-Mailer: git-send-email 2.17.1 In-Reply-To: <20200820210711.8017-1-sandra@codesourcery.com> References: <20200820210711.8017-1-sandra@codesourcery.com> MIME-Version: 1.0 X-ClientProxiedBy: SVR-ORW-MBX-06.mgc.mentorg.com (147.34.90.206) To svr-orw-mbx-03.mgc.mentorg.com (147.34.90.203) X-Spam-Status: No, score=-9.5 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" 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 gcc/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. 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. --- gcc/fortran/ChangeLog.omp | 8 ++++ gcc/fortran/openmp.c | 50 +++++++++++++++++++++- gcc/testsuite/ChangeLog.omp | 12 ++++++ .../gfortran.dg/goacc/combined-directives.f90 | 19 ++++++-- .../goacc/kernels-loop-annotation-18.f95 | 28 ++++++++++++ .../goacc/kernels-loop-annotation-19.f95 | 29 +++++++++++++ .../goacc/private-explicit-kernels-1.f95 | 7 ++- .../goacc/private-predetermined-kernels-1.f95 | 7 ++- 8 files changed, 151 insertions(+), 9 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-18.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-19.f95 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 + 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 + 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 + 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 + 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