From patchwork Fri Nov 6 11:32:02 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Tobias Burnus X-Patchwork-Id: 1395633 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 4CSJC65NyTz9sVK for ; Fri, 6 Nov 2020 22:32:17 +1100 (AEDT) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 446393A53025; Fri, 6 Nov 2020 11:32:12 +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 D68293A53013; Fri, 6 Nov 2020 11:32:08 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org D68293A53013 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Tobias_Burnus@mentor.com IronPort-SDR: 7S/WU2NKQ02gO7AQ8kDFzCdac1gBWOb/vCLytkPP5MmEJNPj0cLpt7J6N2oLbNyRGtpfrq6syn CfFZGLo+d69HzGxalU7UVKqkQ6QgQC6MW3RmEFnnvCde3kAefkQa4BZ0MeN1noFUgYh5U83TLv qL6gmcomKJdGUii26yduuxOLmkd8rb5/DJnhIwkSllpkuYkp6VNbqFQGn7/LAqF4yENwJb63Ue GIQ1pgOna6geNI6ef3OjW8pOhkxbWGYysbBQeklUpQS3EvbudcVX/mZKWXc/SBM0h8NsmDl+uy TSM= X-IronPort-AV: E=Sophos;i="5.77,456,1596528000"; d="diff'?scan'208";a="54828716" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 06 Nov 2020 03:32:07 -0800 IronPort-SDR: nomLCYg0gjvsNgLZL63ic81k69w3oOOyyfADs4KXMYRAGHkQDyqchqtvo/vT7M2xoa/al0hZ83 1xW/OlvmaSN/fvohMwdNdQWMsFF/Z33pFVM5XlyakbO+EAwMUICySjOlrRrnXCxnx/J5Ale2RM JvuAxjd3Kl8N/YVYTmEDDPW2rZFvp7zwCtQqqjYcO9oAT6+Dk84hh+MOpe30vFDes+BhiwGWkF V7NF+d5k7kjqe3uFCyhIaoL2txFOuG/kMhLoJmA4BVRSQkRVsSGI8fJ7gfRXbWmaynduov1kXv 3p8= To: gcc-patches , fortran , Thomas Schwinge From: Tobias Burnus Subject: [Patch, committed] OpenACC/Fortran: Reject '!$acc atomic update capture' Message-ID: <3b1ae8e8-0849-0d6f-599d-d036bd860a95@codesourcery.com> Date: Fri, 6 Nov 2020 12:32:02 +0100 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:78.0) Gecko/20100101 Thunderbird/78.4.0 MIME-Version: 1.0 Content-Language: en-US X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-07.mgc.mentorg.com (139.181.222.7) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-12.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, RCVD_IN_DNSWL_LOW, 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" This patch adds some OpenACC atomic testcases, based on my just committed C/C++ patch. https://gcc.gnu.org/pipermail/gcc-patches/2020-November/558288.html And it fixes an issue I introduced when updating the Fortran atomic handling (for OpenMP and a bit OpenACC) at https://gcc.gnu.org/pipermail/gcc-patches/2020-November/557734.html Namely: For C/C++, OpenACC 2.5 and later has: #pragma acc atomic [atomic-clause] new-line #pragma acc atomic update capture new-line Where atomic-clause is one of read, write, update, or capture. But for Fortran, it only shows for all four clauses "!$acc atomic" followed by the respective clause. Hence, !$acc atomic update capture is not permitted. This patch now removes the support for the latter and adds the testcases. I have also opened the OpenACC spec Issue #333 regarding the C/C++ vs. Fortran difference and the the under-specification of 'update capture'. Committed as r11-4775-gc2e9f586fde57e64dc20e5528870d06cde894785 Tobias ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter commit c2e9f586fde57e64dc20e5528870d06cde894785 Author: Tobias Burnus Date: Fri Nov 6 12:30:20 2020 +0100 OpenACC/Fortran: Reject '!$acc atomic update capture' gcc/fortran/ChangeLog: * openmp.c (gfc_match_oacc_atomic): No longer accept 'update capture'. gcc/testsuite/ChangeLog: * gfortran.dg/goacc-gomp/goacc-gomp.exp: New. * gfortran.dg/goacc-gomp/atomic.f90: New test. * gfortran.dg/goacc/atomic.f90: New test. --- gcc/fortran/openmp.c | 7 +--- gcc/testsuite/gfortran.dg/goacc-gomp/atomic.f90 | 48 ++++++++++++++++++++++ .../gfortran.dg/goacc-gomp/goacc-gomp.exp | 37 +++++++++++++++++ gcc/testsuite/gfortran.dg/goacc/atomic.f90 | 35 ++++++++++++++++ 4 files changed, 122 insertions(+), 5 deletions(-) diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 6cb4f2862ab..1891ac5591b 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -4181,8 +4181,7 @@ gfc_match_omp_atomic (void) } -/* acc atomic [ read | write | update | capture] - acc atomic update capture. */ +/* acc atomic [ read | write | update | capture] */ match gfc_match_oacc_atomic (void) @@ -4191,9 +4190,7 @@ gfc_match_oacc_atomic (void) c->atomic_op = GFC_OMP_ATOMIC_UPDATE; c->memorder = OMP_MEMORDER_RELAXED; gfc_gobble_whitespace (); - if (gfc_match ("update capture") == MATCH_YES) - c->capture = true; - else if (gfc_match ("update") == MATCH_YES) + if (gfc_match ("update") == MATCH_YES) ; else if (gfc_match ("read") == MATCH_YES) c->atomic_op = GFC_OMP_ATOMIC_READ; diff --git a/gcc/testsuite/gfortran.dg/goacc-gomp/atomic.f90 b/gcc/testsuite/gfortran.dg/goacc-gomp/atomic.f90 new file mode 100644 index 00000000000..59186a20982 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc-gomp/atomic.f90 @@ -0,0 +1,48 @@ +! { dg-do compile } */ +! { dg-additional-options "-fdump-tree-original" } */ + +subroutine foo + !$omp requires atomic_default_mem_order(acq_rel) + integer :: i, v + + !$omp atomic read + i = v + + !$acc atomic read + i = v + + !$omp atomic write + i = v + + !$acc atomic write + i = v + + !$omp atomic update + i = i + 1 + + !$acc atomic update + i = i + 1 + + !$omp atomic capture + i = i + 1 + v = i + !$omp end atomic + + !$acc atomic capture + i = i + 1 + v = i + !$acc end atomic + + ! Valid in C/C++ since OpenACC 2.5 but not in Fortran: + ! !$acc atomic update capture + ! i = i + 1 + ! v = i + ! !$acc end atomic +end + +! { dg-final { scan-tree-dump-times "i = #pragma omp atomic read acquire" 1 "original" } } +! { dg-final { scan-tree-dump-times "i = #pragma omp atomic read relaxed" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp atomic release" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp atomic relaxed" 2 "original" } } +! { dg-final { scan-tree-dump-times "v = #pragma omp atomic capture acq_rel" 1 "original" } } +! { dg-final { scan-tree-dump-times "v = #pragma omp atomic capture relaxed" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc-gomp/goacc-gomp.exp b/gcc/testsuite/gfortran.dg/goacc-gomp/goacc-gomp.exp new file mode 100644 index 00000000000..6073fb3978a --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc-gomp/goacc-gomp.exp @@ -0,0 +1,37 @@ +# Copyright (C) 2005-2020 Free Software Foundation, Inc. +# +# This file is part of GCC. +# +# GCC is free software; you can redistribute it and/or modify +# it under the terms of the GNU General Public License as published by +# the Free Software Foundation; either version 3, or (at your option) +# any later version. +# +# GCC is distributed in the hope that it will be useful, +# but WITHOUT ANY WARRANTY; without even the implied warranty of +# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the +# GNU General Public License for more details. +# +# You should have received a copy of the GNU General Public License +# along with GCC; see the file COPYING3. If not see +# . + +# GCC testsuite that uses the `dg.exp' driver. + +# Load support procs. +load_lib gfortran-dg.exp + +if { ![check_effective_target_fopenacc] \ + || ![check_effective_target_fopenmp] } { + return +} + +# Initialize `dg'. +dg-init + +# Main loop. +gfortran-dg-runtest [lsort \ + [find $srcdir/$subdir *.\[fF\]{,90,95,03,08} ] ] "" "-fopenacc -fopenmp" + +# All done. +dg-finish diff --git a/gcc/testsuite/gfortran.dg/goacc/atomic.f90 b/gcc/testsuite/gfortran.dg/goacc/atomic.f90 new file mode 100644 index 00000000000..072d024b132 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/atomic.f90 @@ -0,0 +1,35 @@ +! { dg-do compile } + +subroutine foo + integer :: i, v + !$acc atomic read bar ! { dg-error "21: Unexpected junk after !.ACC ATOMIC statement" } + i = v + + !$acc atomic read write ! { dg-error "21: Unexpected junk after !.ACC ATOMIC statement" } + i = v + + !$acc atomic read seq_cst ! { dg-error "21: Unexpected junk after !.ACC ATOMIC statement" } + i = v + + !$acc atomic read relaxed ! { dg-error "21: Unexpected junk after !.ACC ATOMIC statement" } + i = v + + !$acc atomic update hint(1) ! { dg-error "23: Unexpected junk after !.ACC ATOMIC statement" } + i = i + 1 + + !$acc atomic update update capture ! { dg-error "23: Unexpected junk after !.ACC ATOMIC statement" } + i = i + 1 + v = i + + !$acc atomic update capture capture ! { dg-error "23: Unexpected junk after !.ACC ATOMIC statement" } + i = i + 1 + v = i + + !$acc atomic write capture ! { dg-error "22: Unexpected junk after !.ACC ATOMIC statement" } + i = 1 + + ! Valid in C/C++ since OpenACC 2.5 but not in Fortran: + !$acc atomic update capture ! { dg-error "23: Unexpected junk after !.ACC ATOMIC statement" } + i = i + 1 + v = i +end