diff mbox series

[committed] OpenACC/Fortran: Reject '!$acc atomic update capture'

Message ID 3b1ae8e8-0849-0d6f-599d-d036bd860a95@codesourcery.com
State New
Headers show
Series [committed] OpenACC/Fortran: Reject '!$acc atomic update capture' | expand

Commit Message

Tobias Burnus Nov. 6, 2020, 11:32 a.m. UTC
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
diff mbox series

Patch

commit c2e9f586fde57e64dc20e5528870d06cde894785
Author: Tobias Burnus <tobias@codesourcery.com>
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
+# <http://www.gnu.org/licenses/>.
+
+# 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