From patchwork Wed Dec 15 15:54:15 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Frederik Harwath X-Patchwork-Id: 1568373 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=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Received: from sourceware.org (ip-8-43-85-97.sourceware.org [8.43.85.97]) (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 bilbo.ozlabs.org (Postfix) with ESMTPS id 4JDg235M61z9sXS for ; Thu, 16 Dec 2021 03:00:27 +1100 (AEDT) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id A66BA385783A for ; Wed, 15 Dec 2021 16:00:25 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 528E8385842E; Wed, 15 Dec 2021 15:55:34 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 528E8385842E Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com IronPort-SDR: SWVa975c/qLkZtYxQaOKR5H5ccW/93Rd73gsPNTX9svUe4/rwPtarp5hWlImB+aKWovro/6FXJ 9BVx5qAH9vpBMnruSY9ewYNmy1UDKltJffm3UrSP/pdHroliFVGmB+X0S95cDSSm+5mSnCvubM YVeNse+5emsdkYwmgteNoAjgjlkLSLh+RqY5eViySrgAp1tMbyaNJq5KOBNyAv5dY6gk72TNti UNO1le+3Z/X/W+Y9Fvn1HLUjFlhhjWcidiZSqW8kRIi/4d4ppUtEPiGrljKK26faBP8R/NUaQA uc6P+qKnMtRsUaU1woHcmwDn X-IronPort-AV: E=Sophos;i="5.88,207,1635235200"; d="scan'208";a="69738357" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 15 Dec 2021 07:55:32 -0800 IronPort-SDR: LPQy8uZoKzy2ji5WtEO1jji7bJygHO1TIgnRqmC77ngLMt2T0pwQPEWDS0pteboXKo4oOtmSjS goxW4abRjAekJLjzBG9dBky/ZdwmZU/mO/KATk6vk0x6W9x0C2tgQ9N5csgO/mV7M/BUuazH+L Ba59zteKt6thU0kR0RRoCrRdKHqi4gqcGsnS0CuFk58h8YwX5B3i61HI75ulI2XtgNhxbljM15 uZO4t000zChixVjS3YOpNq/4V9K3dgSCEEwajnk0UTjTVA9PyCj3uCsCgEWzzeds+DoTrfJ8Hw Vv8= From: Frederik Harwath To: Subject: [PATCH 08/40] Annotate inner loops in "acc kernels loop" directives (Fortran). Date: Wed, 15 Dec 2021 16:54:15 +0100 Message-ID: <20211215155447.19379-9-frederik@codesourcery.com> X-Mailer: git-send-email 2.33.0 In-Reply-To: <20211215155447.19379-1-frederik@codesourcery.com> References: <20211215155447.19379-1-frederik@codesourcery.com> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-02.mgc.mentorg.com (139.181.222.2) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-12.6 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.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) 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: , Cc: tobias@codesourcery.com, Sandra Loosemore , thomas@codesourcery.com, fortran@gcc.gnu.org Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" From: Sandra Loosemore 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/openmp.c | 50 ++++++++++++++++++- .../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 ++- 6 files changed, 131 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 -- 2.33.0 ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 243b5e0a9ac6..b0b68b494778 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -9272,7 +9272,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 @@ -9317,6 +9316,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/gfortran.dg/goacc/combined-directives.f90 b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 index 956349204f4d..562a4e40cd7d 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 000000000000..e4e210a92dbb --- /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 000000000000..5dd6e7f538a6 --- /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 5d563d226b0c..0c47045df9c8 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 12a7854526a9..3357a20263e7 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