Background: In real-world code, one can find: !$ACC DECLARE COPYIN(c1es, c2es, ...) as here for the ICON weather model. This clearly implies that other compilers accept and, potentially, require those. For better compatibility with real-world use, the just released OpenACC 3.4 now permits PARAMETER but permits compilers to ignore those (remove them when doing optimizations).
Thus, this patch permits now named constants (PARAMETER) as 'var' in OpenACC [with an off-by-default warning in all but one case (device_resident, no warning)] but then ignores them later. If you look at the following patch, I think the following is ponder about: * Does skipping over PARAMETERS (named constants) in trans-openmp.cc clause handling will break some unrelated OpenACC or OpenMP code? (In principle, resolving an expression should remove the parameter, replacing it by its value. And the called trans-openmp.cc functions also should only deal with non-expressions.) * Does this handle for OpenACC all cases (or did I miss one?) Does it handle too much for OpenACC (or OpenACC?) * Do you think the warning handling is fine/consistent? I think the patch should be fine, but, of course, I might have missed something. Comments, remarks, suggestions about this patch? Tobias
Fortran/OpenACC: Permit PARAMETER as 'var' in clauses (+ ignore) It turned out that other compilers permit (require?) named constants to appear in clauses - and programs actually use this. OpenACC 3.4 added therefore the following: In this spec, a _var_ (in italics) is one of the following: ... * a named constant in Fortran. plus If during an optimization phase _var_ is removed by the compiler, appearances of var in data clauses are ignored. Thus, all errors related to PARAMETER are now downgreaded, most to a -Wsurprising warning, but for 'acc declare device_resident' (which kind of makes sense), no warning is printed. In trans-openmp.cc, those are ignored, unless I missed some code path. (If so, I hope the middle end removes them; but before removing them for the covered cases, the program just compiled & linked fine.) Note that 'ignore PARAMETER inside clauses' in trans-openmp.cc would in principle also apply to expressions ('if (var)') but those should be evaluated during 'resolve.cc' + 'openmp.cc' to their (numeric, logical, string) value such that there should be no issue. gcc/fortran/ChangeLog: * invoke.texi (-Wsurprising): Note about OpenACC warning related to PARAMATER. * openmp.cc (resolve_omp_clauses, gfc_resolve_oacc_declare): Accept PARAMETER for OpenACC but add surprising warning. * trans-openmp.cc (gfc_trans_omp_variable_list, gfc_trans_omp_clauses): Ignore PARAMETER inside clauses. gcc/testsuite/ChangeLog: * gfortran.dg/goacc/parameter.f95: Add -Wsurprising flag and update expected diagnostic. * gfortran.dg/goacc/parameter-3.f90: New test. * gfortran.dg/goacc/parameter-4.f90: New test. gcc/fortran/invoke.texi | 4 ++++ gcc/fortran/openmp.cc | 30 ++++++++++++++++++++----- gcc/fortran/trans-openmp.cc | 13 ++++++++--- gcc/testsuite/gfortran.dg/goacc/parameter-3.f90 | 16 +++++++++++++ gcc/testsuite/gfortran.dg/goacc/parameter-4.f90 | 26 +++++++++++++++++++++ gcc/testsuite/gfortran.dg/goacc/parameter.f95 | 27 +++++++++++----------- 6 files changed, 94 insertions(+), 22 deletions(-) diff --git a/gcc/fortran/invoke.texi b/gcc/fortran/invoke.texi index da085d124f9..0b893e876a5 100644 --- a/gcc/fortran/invoke.texi +++ b/gcc/fortran/invoke.texi @@ -1170,6 +1170,10 @@ A @code{CHARACTER} variable is declared with negative length. With @option{-fopenmp}, for fixed-form source code, when an @code{omx} vendor-extension sentinel is encountered. (The equivalent @code{ompx}, used in free-form source code, is diagnosed by default.) + +@item +With @option{-fopenacc}, when using named constances with clauses that +take a variable as doing so has no effect. @end itemize @opindex Wtabs diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index fe0a47a6948..f1acc00f561 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -8895,15 +8895,21 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, if (list == OMP_LIST_MAP && n->sym->attr.flavor == FL_PARAMETER) { + /* OpenACC since 3.4 permits for Fortran named constants, but + permits removing then as optimization is not needed and such + ignore them. Likewise below for FIRSTPRIVATE. */ if (openacc) - gfc_error ("Object %qs is not a variable at %L; parameters" - " cannot be and need not be copied", n->sym->name, - &n->where); + gfc_warning (OPT_Wsurprising, "Clause for object %qs at %L is " + "ignored as parameters need not be copied", + n->sym->name, &n->where); else gfc_error ("Object %qs is not a variable at %L; parameters" " cannot be and need not be mapped", n->sym->name, &n->where); } + else if (openacc && n->sym->attr.flavor == FL_PARAMETER) + gfc_warning (OPT_Wsurprising, "Clause for object %qs at %L is ignored" + " as it is a parameter", n->sym->name, &n->where); else if (list != OMP_LIST_USES_ALLOCATORS) gfc_error ("Object %qs is not a variable at %L", n->sym->name, &n->where); @@ -12756,9 +12762,21 @@ gfc_resolve_oacc_declare (gfc_namespace *ns) && (n->sym->attr.flavor != FL_PROCEDURE || n->sym->result != n->sym)) { - gfc_error ("Object %qs is not a variable at %L", - n->sym->name, &oc->loc); - continue; + if (n->sym->attr.flavor != FL_PARAMETER) + { + gfc_error ("Object %qs is not a variable at %L", + n->sym->name, &oc->loc); + continue; + } + /* Note that OpenACC 3.4 permits name constants, but the + implementation is permitted to ignore the clause; + as semantically, device_resident kind of makes sense + (and the wording with it is a bit odd), the warning + is suppressed. */ + if (list != OMP_LIST_DEVICE_RESIDENT) + gfc_warning (OPT_Wsurprising, "Object %qs at %L is ignored as" + " parameters need not be copied", n->sym->name, + &oc->loc); } if (n->expr && n->expr->ref->type == REF_ARRAY) diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index a2e70fca0b3..f3d7cd4ffee 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -2792,8 +2792,13 @@ gfc_trans_omp_variable_list (enum omp_clause_code code, gfc_omp_namelist *namelist, tree list, bool declare_simd) { + /* PARAMETER (named constants) are excluded as OpenACC 3.4 permits them now + as 'var' but permits compilers to ignore them. In expressions, it should + have been replaced by the value (and this function should not be called + anyway) and for var-using clauses, they should just be skipped. */ for (; namelist != NULL; namelist = namelist->next) - if (namelist->sym->attr.referenced || declare_simd) + if ((namelist->sym->attr.referenced || declare_simd) + && namelist->sym->attr.flavor != FL_PARAMETER) { tree t = gfc_trans_omp_variable (namelist->sym, declare_simd); if (t != error_mark_node) @@ -4029,7 +4034,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, case OMP_LIST_MAP: for (; n != NULL; n = n->next) { - if (!n->sym->attr.referenced) + if (!n->sym->attr.referenced + || n->sym->attr.flavor == FL_PARAMETER) continue; location_t map_loc = gfc_get_location (&n->where); @@ -4986,7 +4992,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, case OMP_LIST_CACHE: for (; n != NULL; n = n->next) { - if (!n->sym->attr.referenced) + if (!n->sym->attr.referenced + && n->sym->attr.flavor != FL_PARAMETER) continue; switch (list) diff --git a/gcc/testsuite/gfortran.dg/goacc/parameter-3.f90 b/gcc/testsuite/gfortran.dg/goacc/parameter-3.f90 new file mode 100644 index 00000000000..2c8aa618cc2 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/parameter-3.f90 @@ -0,0 +1,16 @@ +! { dg-do compile } + +subroutine x + integer :: var + integer, parameter :: ilog = 0 + integer, parameter :: array(*) = [11,22,33] + !$ACC DECLARE COPYIN(ilog, array, var, array) ! { dg-error "Symbol 'array' present on multiple clauses" } +end subroutine x + +integer :: a +integer, parameter :: b = 4 +integer, parameter :: c(*) = [1,2,3] + +!$acc parallel copy(a,c,b,c) ! { dg-error "Symbol 'c' present on multiple clauses" } +!$acc end parallel +end diff --git a/gcc/testsuite/gfortran.dg/goacc/parameter-4.f90 b/gcc/testsuite/gfortran.dg/goacc/parameter-4.f90 new file mode 100644 index 00000000000..aadd7cffe81 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/parameter-4.f90 @@ -0,0 +1,26 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-original" } + +subroutine x + integer :: var + integer, parameter :: ilog = 0 + integer, parameter :: array(*) = [11,22,33] + !$ACC DECLARE COPYIN(ilog, array, var) +end subroutine x + +integer :: a +integer, parameter :: b = 4 +integer, parameter :: c(*) = [1,2,3] + +!$acc parallel copy(a,c,b) + a = c(2) + b +!$acc end parallel + +!$acc parallel firstprivate(a,c,b) + a = c(2) + b +!$acc end parallel +end + +! { dg-final { scan-tree-dump-times "#pragma acc data map\\(to:var\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(tofrom:a\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma acc parallel firstprivate\\(a\\)" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/parameter.f95 b/gcc/testsuite/gfortran.dg/goacc/parameter.f95 index b5813386721..a9bde4a67ec 100644 --- a/gcc/testsuite/gfortran.dg/goacc/parameter.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/parameter.f95 @@ -1,4 +1,5 @@ ! { dg-do compile } +! { dg-additional-options "-Wsurprising" } module test contains @@ -6,37 +7,37 @@ contains implicit none integer :: i integer, parameter :: a = 1 - !$acc declare device_resident (a) ! { dg-error "is not a variable" } - !$acc data copy (a) ! { dg-error "not a variable" } + !$acc declare device_resident (a) ! (no warning here - for semi-good reasons) + !$acc data copy (a) ! { dg-warning "Clause for object 'a' at .1. is ignored as parameters need not be copied \\\[-Wsurprising\\\]" } !$acc end data - !$acc data deviceptr (a) ! { dg-error "not a variable" } + !$acc data deviceptr (a) ! { dg-warning "Clause for object 'a' at .1. is ignored as parameters need not be copied \\\[-Wsurprising\\\]" } !$acc end data - !$acc parallel private (a) ! { dg-error "not a variable" } + !$acc parallel private (a) ! { dg-warning "Clause for object 'a' at .1. is ignored as it is a parameter \\\[-Wsurprising\\\]" } !$acc end parallel - !$acc serial private (a) ! { dg-error "not a variable" } + !$acc serial private (a) ! { dg-warning "Clause for object 'a' at .1. is ignored as it is a parameter \\\[-Wsurprising\\\]" } !$acc end serial - !$acc host_data use_device (a) ! { dg-error "not a variable" } + !$acc host_data use_device (a) ! { dg-warning "Clause for object 'a' at .1. is ignored as it is a parameter \\\[-Wsurprising\\\]" } !$acc end host_data - !$acc parallel loop reduction(+:a) ! { dg-error "not a variable" } + !$acc parallel loop reduction(+:a) ! { dg-warning "Clause for object 'a' at .1. is ignored as it is a parameter \\\[-Wsurprising\\\]" } do i = 1,5 enddo !$acc end parallel loop - !$acc serial loop reduction(+:a) ! { dg-error "not a variable" } + !$acc serial loop reduction(+:a) ! { dg-warning "Clause for object 'a' at .1. is ignored as it is a parameter \\\[-Wsurprising\\\]" } do i = 1,5 enddo !$acc end serial loop !$acc parallel loop do i = 1,5 - !$acc cache (a) ! { dg-error "not a variable" } + !$acc cache (a) ! { dg-warning "Clause for object 'a' at .1. is ignored as it is a parameter \\\[-Wsurprising\\\]" } enddo !$acc end parallel loop !$acc serial loop do i = 1,5 - !$acc cache (a) ! { dg-error "not a variable" } + !$acc cache (a) ! { dg-warning "Clause for object 'a' at .1. is ignored as it is a parameter \\\[-Wsurprising\\\]" } enddo !$acc end serial loop - !$acc update device (a) ! { dg-error "not a variable" } - !$acc update host (a) ! { dg-error "not a variable" } - !$acc update self (a) ! { dg-error "not a variable" } + !$acc update device (a) ! { dg-warning "Clause for object 'a' at .1. is ignored as parameters need not be copied \\\[-Wsurprising\\\]" } + !$acc update host (a) ! { dg-warning "Clause for object 'a' at .1. is ignored as parameters need not be copied \\\[-Wsurprising\\\]" } + !$acc update self (a) ! { dg-warning "Clause for object 'a' at .1. is ignored as parameters need not be copied \\\[-Wsurprising\\\]" } end subroutine oacc1 end module test