From patchwork Tue Jul 7 08:52:08 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Frederik Harwath X-Patchwork-Id: 1324202 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@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 [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 ozlabs.org (Postfix) with ESMTPS id 4B1GQv5NPSz9sRK for ; Tue, 7 Jul 2020 18:52:22 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 8EC1A3851C2A; Tue, 7 Jul 2020 08:52:19 +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 109013851C20 for ; Tue, 7 Jul 2020 08:52:15 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 109013851C20 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Frederik_Harwath@mentor.com IronPort-SDR: Nx8DxbeJ7OprBebEA2uXgtI7LPp2s4Ni6xg1G8Bi2uymscS3QZU7kazeQnVIjOn26wVt2esObg ugUZfhJgiOml1ueifr03rePi/ciRmFjwa2qqGkAmt7FdJ4ljtAEsHqs5QMkyQABGxvBbHqFb3s JNTGhVXGFtRlGbLTMutFjo2ldx4gRfrwHRdbHDWQoitIndN8QUJcFRNWmcJjcwNJTh0VzCmdgt cJ54qH0wslrf6skfdBgoFOzOLkQgcJE8xk/fbc2bnpr0N4mdqVEflP123rbWvjh6JVnfpsX5Ri Eks= X-IronPort-AV: E=Sophos;i="5.75,323,1589270400"; d="scan'208,223";a="50662159" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 07 Jul 2020 00:52:14 -0800 IronPort-SDR: p1VL2Fwj6Tag8CnxBjMYyE4uiMdP39tHk/Al1C/60+/mUXthanmspYFZELOyz9SUmr3EjrKTsw YoZEQojdQZzkOyKfexZugowJBdQ6QCFQS4+l6xjzH5NG84Xw2c403e0nUiRkOZs2dMnYRebFGa dRFOz4uTGRJUtzZjM+S9Cy37+30QtkPn1b8nKme5QIcocS0TT7tQE05WuCTKmFpjY/s1ulwose NTwGbvI7smdGU3GyP9Gm7qpb90iu+ORWpIVCqNcDQKemfopSjLOxQTbbDCFDXGI+nsFM+uvz3i yUQ= From: Frederik Harwath To: , Thomas Schwinge Subject: [PATCH] [og10] libgomp, Fortran: Fix OpenACC "gang reduction on an orphan loop" error message Date: Tue, 7 Jul 2020 10:52:08 +0200 Message-ID: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) To SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) X-Spam-Status: No, score=-13.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: , Cc: Jakub Jelinek , Tobias Burnus Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" Hi, This patch fixes the check for reductions on orphaned gang loops in the Fortran frontend which (in contrast to the C, C++ frontends) erroneously rejects reductions on gang loops that are contained in "kernels" constructs and which hence are not orphaned. According to the OpenACC standard version 2.5 and later, reductions on orphaned gang loops are explicitly disallowed (cf. section "Changes from Version 2.0 to 2.5"). Remember that a loop is "orphaned" if it is not lexically contained in a compute construct (cf. section "Loop construct" of the OpenACC standard), i.e. in either a "parallel", a "serial", or a "kernels" construct. The patch has been tested by running the GCC and libgomp testsuites. The latter tests ran with offloading to nvptx although that should not be important here unless there was some very subtle reason for forbidding the gang reductions on kernels loops. As expect, there seems to be no such reason, i.e. I observed no regressions with the patch. Can I include the patch in OG10? Best regards, Frederik ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter From 7320635211fff3a773beb0de1914dbfcc317ab37 Mon Sep 17 00:00:00 2001 From: Frederik Harwath Date: Tue, 7 Jul 2020 10:41:21 +0200 Subject: [PATCH] libgomp, Fortran: Fix OpenACC "gang reduction on an orphan loop" error message According to the OpenACC standard version 2.5 and later, reductions on orphaned gang loops are explicitly disallowed (cf. section "Changes from Version 2.0 to 2.5"). A loop is "orphaned" if it is not lexically contained in a compute construct (cf. section "Loop construct" of the OpenACC standard), i.e. in either a "parallel", a "serial", or a "kernels" construct. This commit fixes the check for reductions on orphaned gang loops in the Fortran frontend which (in contrast to the C, C++ frontends) erroneously rejects reductions on gang loops that are contained in "kernels" constructs. 2020-07-07 Frederik Harwath gcc/fortran/ * openmp.c (oacc_is_parallel_or_serial): Removed function. (oacc_is_kernels): New function. (oacc_is_compute_construct): New function. (resolve_oacc_loop_blocks): Use "oacc_is_compute_construct" instead of "oacc_is_parallel_or_serial" for checking that a loop is not orphaned. gcc/testsuite/ * gfortran.dg/goacc/orphan-reductions-2.f90: New test verifying that the error message is not emitted for non-orphaned loops. * c-c++-common/goacc/orphan-reductions-2.c: Likewise for C and C++. --- gcc/fortran/openmp.c | 13 +++- .../c-c++-common/goacc/orphan-reductions-2.c | 69 +++++++++++++++++++ .../gfortran.dg/goacc/orphan-reductions-2.f90 | 58 ++++++++++++++++ 3 files changed, 137 insertions(+), 3 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c create mode 100644 gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90 diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 28408c4c99a..83c498112a8 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -5926,9 +5926,16 @@ oacc_is_serial (gfc_code *code) } static bool -oacc_is_parallel_or_serial (gfc_code *code) +oacc_is_kernels (gfc_code *code) { - return oacc_is_parallel (code) || oacc_is_serial (code); + return code->op == EXEC_OACC_KERNELS || code->op == EXEC_OACC_KERNELS_LOOP; +} + +static bool +oacc_is_compute_construct (gfc_code *code) +{ + return oacc_is_parallel (code) || oacc_is_serial (code) + || oacc_is_kernels (code); } static gfc_statement @@ -6222,7 +6229,7 @@ resolve_oacc_loop_blocks (gfc_code *code) for (c = omp_current_ctx; c; c = c->previous) if (!oacc_is_loop (c->code)) break; - if (c == NULL || !oacc_is_parallel_or_serial (c->code)) + if (c == NULL || !oacc_is_compute_construct (c->code)) gfc_error ("gang reduction on an orphan loop at %L", &code->loc); } diff --git a/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c new file mode 100644 index 00000000000..2b651fd2b9f --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-2.c @@ -0,0 +1,69 @@ +/* Verify that the error message for gang reduction on orphaned OpenACC loops + is not reported for non-orphaned loops. */ + +#include + +int +kernels (int n) +{ + int i, s1 = 0, s2 = 0; +#pragma acc kernels + { +#pragma acc loop gang reduction(+:s1) /* { dg-bogus "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s1 = s1 + 2; + +#pragma acc loop gang reduction(+:s2) /* { dg-bogus "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s2 = s2 + 2; + } + return s1 + s2; +} + +int +parallel (int n) +{ + int i, s1 = 0, s2 = 0; +#pragma acc parallel + { +#pragma acc loop gang reduction(+:s1) /* { dg-bogus "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s1 = s1 + 2; + +#pragma acc loop gang reduction(+:s2) /* { dg-bogus "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s2 = s2 + 2; + } + return s1 + s2; +} + +int +parallel_combined (int n) +{ + int i, s1 = 0, s2 = 0; +#pragma acc parallel loop gang reduction(+:s1) /* { dg-bogus "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s1 = s1 + 2; + +#pragma acc parallel loop gang reduction(+:s2) /* { dg-bogus "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s2 = s2 + 2; + + return s1 + s2; +} + +int +kernels_combined (int n) +{ + int i, s1 = 0, s2 = 0; +#pragma acc kernels loop gang reduction(+:s1) /* { dg-bogus "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s1 = s1 + 2; + +#pragma acc kernels loop gang reduction(+:s2) /* { dg-bogus "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s2 = s2 + 2; + + return s1 + s2; +} + diff --git a/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90 b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90 new file mode 100644 index 00000000000..13887a059fe --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-2.f90 @@ -0,0 +1,58 @@ +! Verify that the error message for gang reductions on orphaned OpenACC loops +! is not reported for non-orphaned loops. + +subroutine kernels + implicit none + + integer, parameter :: n = 100 + integer :: i, sum + sum = 0 + + !$acc kernels + !$acc loop gang reduction(+:sum) ! { dg-bogus "gang reduction on an orphan loop" } + do i = 1, n + sum = sum + 1 + end do + !$acc end kernels +end subroutine kernels + +subroutine parallel + implicit none + + integer, parameter :: n = 100 + integer :: i, sum + sum = 0 + + !$acc parallel + !$acc loop gang reduction(+:sum) ! { dg-bogus "gang reduction on an orphan loop" } + do i = 1, n + sum = sum + 1 + end do + !$acc end parallel +end subroutine parallel + +subroutine kernels_combined + implicit none + + integer, parameter :: n = 100 + integer :: i, sum + sum = 0 + + !$acc kernels loop gang reduction(+:sum) ! { dg-bogus "gang reduction on an orphan loop" } + do i = 1, n + sum = sum + 1 + end do +end subroutine kernels_combined + +subroutine parallel_combined + implicit none + + integer, parameter :: n = 100 + integer :: i, sum + sum = 0 + + !$acc parallel loop gang reduction(+:sum) ! { dg-bogus "gang reduction on an orphan loop" } + do i = 1, n + sum = sum + 1 + end do +end subroutine parallel_combined -- 2.17.1