From patchwork Wed Dec 15 15:54:14 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Frederik Harwath X-Patchwork-Id: 1568372 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+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) 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 bilbo.ozlabs.org (Postfix) with ESMTPS id 4JDg1Z6vLpz9sXS for ; Thu, 16 Dec 2021 03:00:02 +1100 (AEDT) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id D7F8D385782A for ; Wed, 15 Dec 2021 16:00:00 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id E9975385802B for ; Wed, 15 Dec 2021 15:55:23 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org E9975385802B 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: BlF9C89zElbXZJpSbHatRGltKk0zXjkliH0LRjwYz8yYSmHQDfBA36b8zxudMGKK5NRf7eqONI dK2UQP8ljYMaEk4eNSlP6qAnnyfxe13HfgBeQXoA2A+XloL0MSWF4RveW+Cba/Yfp19kVJr8JF hzjIjoZpKl7L1NSwu+E/GTps4TnPb7wWEqGYxZkuanpQJmjZNvJJZxX28azbOx6wHf8QyepPfv xrRb9GpGfF7vEaWyTZQ6hIzKClLRCnlRya0ar+5bEe4hQ4dbAEQbBzQZGpaqmZS1F9yChZyxGn 129wDyOdmO2WGaF6WzZTv6xM X-IronPort-AV: E=Sophos;i="5.88,207,1635235200"; d="scan'208";a="69736543" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 15 Dec 2021 07:55:22 -0800 IronPort-SDR: 3whl+JL5uVbHSQwlU4P0ct6jFfbFVwPSD9MYSYBeqLUNP6Htc+b59lwZVJ5zksQfnq6A5nzZS2 6mQuBWHrfatuIZEeEb23MGQMu+TD6kHFbDrztdqqzob1I54+9Y+ZfhiLcRu4SZ8cxXryb5rSkt N03chzgs97iIaU2JB4Mu1fNZjJaO+FJ10rPB/10sI9zAxmdF0bfDw7JYlaCFcAbLnoNyuBosPh FHlXUWhJiWCsFthr3Ef4w9EuIIEB7ENgq+SLcarwuThlkd5302iTql3hAEOakG8bmmzVVPE+dZ 7GA= From: Frederik Harwath To: Subject: [PATCH 07/40] Annotate inner loops in "acc kernels loop" directives (C/C++). Date: Wed, 15 Dec 2021 16:54:14 +0100 Message-ID: <20211215155447.19379-8-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-10.mgc.mentorg.com (139.181.222.10) 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: Sandra Loosemore , thomas@codesourcery.com, joseph@codesourcery.com, nathan@acm.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 for C and C++. 2020-08-19 Sandra Loosemore gcc/c-family/ * c-omp.c (annotate_loops_in_kernels_regions): Process inner loops in combined "acc kernels loop" directives. gcc/testsuite/ * c-c++-common/goacc/kernels-loop-annotation-18.c: New. * c-c++-common/goacc/kernels-loop-annotation-19.c: New. * c-c++-common/goacc/combined-directives.c: Adjust expected patterns. --- gcc/c-family/c-omp.c | 36 ++++++++++++------- .../c-c++-common/goacc/combined-directives.c | 2 +- .../goacc/kernels-loop-annotation-18.c | 18 ++++++++++ .../goacc/kernels-loop-annotation-19.c | 19 ++++++++++ 4 files changed, 62 insertions(+), 13 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-18.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-19.c -- 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/c-family/c-omp.c b/gcc/c-family/c-omp.c index fad50da8fbc4..30757877eafe 100644 --- a/gcc/c-family/c-omp.c +++ b/gcc/c-family/c-omp.c @@ -3477,18 +3477,30 @@ annotate_loops_in_kernels_regions (tree *nodeptr, int *walk_subtrees, /* Do not try to add automatic OpenACC annotations inside manually annotated loops. Presumably, the user avoided doing it on purpose; for example, all available levels of parallelism may - have been used up. */ - { - struct annotation_info nested_info - = { NULL_TREE, NULL_TREE, false, as_explicit_annotation, - node, info }; - if (info->state >= as_in_kernels_region) - do_not_annotate_loop_nest (info, as_explicit_annotation, - node); - walk_tree (&OMP_BODY (node), annotate_loops_in_kernels_regions, - (void *) &nested_info, NULL); - *walk_subtrees = 0; - } + have been used up. However, assume that the combined construct + "#pragma acc kernels loop" means to try to process the whole + loop nest. + Note that a single OACC_LOOP construct represents an entire set + of collapsed loops so we do not have to deal explicitly with the + collapse clause here, as the Fortran front end does. */ + if (info->state == as_in_kernels_region && OACC_LOOP_COMBINED (node)) + { + walk_tree (&OMP_BODY (node), annotate_loops_in_kernels_regions, + (void *) info, NULL); + *walk_subtrees = 0; + } + else + { + struct annotation_info nested_info + = { NULL_TREE, NULL_TREE, false, as_explicit_annotation, + node, info }; + if (info->state >= as_in_kernels_region) + do_not_annotate_loop_nest (info, as_explicit_annotation, + node); + walk_tree (&OMP_BODY (node), annotate_loops_in_kernels_regions, + (void *) &nested_info, NULL); + *walk_subtrees = 0; + } break; case FOR_STMT: diff --git a/gcc/testsuite/c-c++-common/goacc/combined-directives.c b/gcc/testsuite/c-c++-common/goacc/combined-directives.c index c2a3c57b48b8..2519f23d49f0 100644 --- a/gcc/testsuite/c-c++-common/goacc/combined-directives.c +++ b/gcc/testsuite/c-c++-common/goacc/combined-directives.c @@ -110,7 +110,7 @@ test () // { dg-final { scan-tree-dump-times "acc loop worker" 2 "gimple" } } // { dg-final { scan-tree-dump-times "acc loop vector" 2 "gimple" } } // { dg-final { scan-tree-dump-times "acc loop seq" 2 "gimple" } } -// { dg-final { scan-tree-dump-times "acc loop auto" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "acc loop auto" 6 "gimple" } } // { dg-final { scan-tree-dump-times "acc loop tile.2, 3" 2 "gimple" } } // { dg-final { scan-tree-dump-times "acc loop independent private.i" 2 "gimple" } } // { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-18.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-18.c new file mode 100644 index 000000000000..89ec6447625f --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-18.c @@ -0,0 +1,18 @@ +/* { 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. */ + +void f (float *a, float *b) +{ +#pragma acc kernels loop + for (int k = 0; k < 20; k++) + for (int l = 0; l < 20; l++) + for (int m = 0; m < 20; m++) + b[m] = a[m]; +} + +/* { dg-final { scan-tree-dump-times "acc loop auto" 2 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-19.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-19.c new file mode 100644 index 000000000000..77a3b7a9136d --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-19.c @@ -0,0 +1,19 @@ +/* { 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. */ + +void f (float *a, float *b) +{ +#pragma acc kernels loop collapse(2) + for (int k = 0; k < 20; k++) + for (int l = 0; l < 20; l++) + for (int m = 0; m < 20; m++) + b[m] = a[m]; +} + +/* { dg-final { scan-tree-dump-times "acc loop collapse.2." 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "acc loop auto" 1 "original" } } */