[This fixes a bunch of issues found when actually only wanting to add a check for the following restriction.]
OpenACC's "Declare Directive" has the following restriction: "A declare directive must be in the same scope as the declaration of any var that appears in the data clauses of the directive." The gfortran now rejects a "var" declared is in a different namespace than the "$acc declare". (Use-associated variables are already rejected.) Testing showed that a straight-forward check fails if the result variable is the function name – as then the function symbol is in the parent scope. — Extending the failing test to use a result variable showed that the current is-a-module-variable check didn't work and when fixing, one was running into a likewise issue. The reason that I exclude 's' being a module is that at resolution time, an is-variable check is done. The other changes are because the following restriction was mishandled: "In a Fortran module declaration section, only create, copyin, device_resident, and link clauses are allowed." But all examples where for variables using those in a module procedure … OK for the trunk? Cheers, Tobias PS: For those who wounder what happens in a BLOCK DATA construct: gfortran outrightly rejects 'acc declare'. (It probably should work for COMMON blocks with 'declare device_resident' – but currently the spec only permits it in program/subroutine/function + declaration part of a module.) PPS: The PR shows (for C) that one can construct a test case, which violates the OpenACC restriction and actually fails with an ICE. I have a draft patch for C (see PR) but not yet one for C++, hence, I start with the Fortran side. – I currently still struggle to write a same-scope check in C++. [The C test case in turn was a fallout of debugging an ICE-on-valid-code issue …] ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
[Fortran, OpenACC] Reject vars of different scope in $acc declare (PR94120) 2020-10-03 Tobias Burnus <tob...@codesourcery.com> PR middle-end/94120 * openmp.c (gfc_match_oacc_declare): Accept function-result variables; reject variables declared in a different scoping unit. 2020-10-03 Tobias Burnus <tob...@codesourcery.com> PR middle-end/94120 * gfortran.dg/goacc/pr78260-2.f90: Correct scan-tree-dump-times. Extend test case to result variables. * gfortran.dg/goacc/declare-2.f95: Actually check module-declaration restriction of OpenACC. * gfortran.dg/goacc/declare-3.f95: Remove case where this restriction is violated. * gfortran.dg/goacc/pr94120-1.f90: New. * gfortran.dg/goacc/pr94120-2.f90: New. * gfortran.dg/goacc/pr94120-3.f90: New. gcc/fortran/openmp.c | 12 +++++++++++- gcc/testsuite/gfortran.dg/goacc/declare-2.f95 | 21 ++++++++++++++++----- gcc/testsuite/gfortran.dg/goacc/declare-3.f95 | 10 +--------- gcc/testsuite/gfortran.dg/goacc/pr78260-2.f90 | 13 +++++++++++-- gcc/testsuite/gfortran.dg/goacc/pr94120-1.f90 | 11 +++++++++++ gcc/testsuite/gfortran.dg/goacc/pr94120-2.f90 | 12 ++++++++++++ gcc/testsuite/gfortran.dg/goacc/pr94120-3.f90 | 13 +++++++++++++ 7 files changed, 75 insertions(+), 17 deletions(-) diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 35f6b2f4938..930bca541b9 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -2155,7 +2155,8 @@ gfc_match_oacc_declare (void) { gfc_symbol *s = n->sym; - if (s->ns->proc_name && s->ns->proc_name->attr.proc == PROC_MODULE) + if (gfc_current_ns->proc_name + && gfc_current_ns->proc_name->attr.flavor == FL_MODULE) { if (n->u.map_op != OMP_MAP_ALLOC && n->u.map_op != OMP_MAP_TO) { @@ -2174,6 +2175,15 @@ gfc_match_oacc_declare (void) return MATCH_ERROR; } + if ((s->result == s && s->ns->contained != gfc_current_ns) + || ((s->attr.flavor == FL_UNKNOWN || s->attr.flavor == FL_VARIABLE) + && s->ns != gfc_current_ns)) + { + gfc_error ("Variable %qs shall be declared in the same scoping unit " + "as !$ACC DECLARE at %L", s->name, &where); + return MATCH_ERROR; + } + if ((s->attr.dimension || s->attr.codimension) && s->attr.dummy && s->as->type != AS_EXPLICIT) { diff --git a/gcc/testsuite/gfortran.dg/goacc/declare-2.f95 b/gcc/testsuite/gfortran.dg/goacc/declare-2.f95 index 7aa3dab4707..bad5de9d757 100644 --- a/gcc/testsuite/gfortran.dg/goacc/declare-2.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/declare-2.f95 @@ -1,9 +1,5 @@ module amod - -contains - -subroutine asubr (b) implicit none integer :: b(8) @@ -16,9 +12,24 @@ subroutine asubr (b) !$acc declare present_or_create (b) ! { dg-error "present on multiple" } !$acc declare deviceptr (b) ! { dg-error "Invalid clause in module" } !$acc declare create (b) copyin (b) ! { dg-error "present on multiple" } +end module +module amod2 +contains +subroutine asubr (a, b, c, d, e, f, g, h, i, j, k) + implicit none + integer, dimension(8) :: a, b, c, d, e, f, g, h, i, j, k + + !$acc declare copy (a) + !$acc declare copyout (b) + !$acc declare present (c) + !$acc declare present_or_copy (d) + !$acc declare present_or_copyin (e) + !$acc declare present_or_copyout (f) + !$acc declare present_or_create (g) + !$acc declare deviceptr (h) + !$acc declare create (j) copyin (k) end subroutine - end module module bmod diff --git a/gcc/testsuite/gfortran.dg/goacc/declare-3.f95 b/gcc/testsuite/gfortran.dg/goacc/declare-3.f95 index 80d9903a9dc..9127cba6600 100644 --- a/gcc/testsuite/gfortran.dg/goacc/declare-3.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/declare-3.f95 @@ -14,12 +14,6 @@ module mod_b !$acc declare copyin (b) end module -module mod_c - implicit none - integer :: c - !$acc declare deviceptr (c) -end module - module mod_d implicit none integer :: d @@ -35,7 +29,6 @@ end module subroutine sub1 use mod_a use mod_b - use mod_c use mod_d use mod_e end subroutine sub1 @@ -43,11 +36,10 @@ end subroutine sub1 program test use mod_a use mod_b - use mod_c use mod_d use mod_e - ! { dg-final { scan-tree-dump {(?n)#pragma acc data map\(force_alloc:d\) map\(force_deviceptr:c\) map\(force_to:b\) map\(force_alloc:a\)$} original } } + ! { dg-final { scan-tree-dump {(?n)#pragma acc data map\(force_alloc:d\) map\(force_to:b\) map\(force_alloc:a\)$} original } } end program test ! { dg-final { scan-tree-dump-times {#pragma acc data} 1 original } } diff --git a/gcc/testsuite/gfortran.dg/goacc/pr78260-2.f90 b/gcc/testsuite/gfortran.dg/goacc/pr78260-2.f90 index e28564d6f70..f8a3dc864cc 100644 --- a/gcc/testsuite/gfortran.dg/goacc/pr78260-2.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/pr78260-2.f90 @@ -4,6 +4,8 @@ ! PR fortran/78260 +! Loosely related to PR fortran/94120 + module m implicit none integer :: n = 0 @@ -14,7 +16,14 @@ contains f1 = 5 !$acc end kernels end function f1 + integer function g1() result(g1res) + !$acc declare present(g1res) + !$acc kernels copyin(g1res) + g1res = 5 + !$acc end kernels + end function g1 end module m ! { dg-final { scan-tree-dump-times "#pragma acc data map\\(force_present:__result_f1\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "#pragma acc data map\\(force_present:__result_f1\\)" 1 "original" } } - +! { dg-final { scan-tree-dump-times "#pragma acc kernels map\\(to:__result_f1\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma acc data map\\(force_present:g1res\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma acc kernels map\\(to:g1res\\)" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/pr94120-1.f90 b/gcc/testsuite/gfortran.dg/goacc/pr94120-1.f90 new file mode 100644 index 00000000000..8d1fc16a5ff --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/pr94120-1.f90 @@ -0,0 +1,11 @@ +! { dg-do compile } +! +! PR fortran/94120 +! +implicit none +integer :: i +contains + subroutine f() + !$acc declare copy(i) ! { dg-error "Variable 'i' shall be declared in the same scoping unit as !.ACC DECLARE" } + end +end diff --git a/gcc/testsuite/gfortran.dg/goacc/pr94120-2.f90 b/gcc/testsuite/gfortran.dg/goacc/pr94120-2.f90 new file mode 100644 index 00000000000..216c04b844d --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/pr94120-2.f90 @@ -0,0 +1,12 @@ +! { dg-do compile } +! +! PR fortran/94120 +! +! BLOCK is not supported in OpenACC <= 3.0 +! +subroutine f() + block + integer :: k + !$acc declare copy(j) ! { dg-error "Sorry, !.ACC DECLARE at .1. is not allowed in BLOCK construct" } + end block +end diff --git a/gcc/testsuite/gfortran.dg/goacc/pr94120-3.f90 b/gcc/testsuite/gfortran.dg/goacc/pr94120-3.f90 new file mode 100644 index 00000000000..1eec90ae7ed --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/pr94120-3.f90 @@ -0,0 +1,13 @@ +! { dg-do compile } +! +! PR fortran/94120 +! +! Note: BLOCK is not supported in OpenACC <= 3.0 – but the following check comes earlier: +! It is also invalid because the variable is in a different scoping unit +! +subroutine g() + integer :: k + block + !$acc declare copy(k) ! { dg-error "Variable 'k' shall be declared in the same scoping unit as !.ACC DECLARE" } + end block +end