| Message ID | 20200104022514.115942-2-julian@codesourcery.com |
|---|---|
| State | New |
| Headers | show |
| Series | [1/3] Add OpenACC test for sub-references being pointer or allocatable variables | expand |
Hi Julian, @Jakub: I suggest to enable it also for OpenMP as prep work. Any objections/comments? @Thomas: This all assumes that "copy(a,a.comp)" is not valid in OpenACC; does this match your understanding of the spec? On 1/4/20 3:25 AM, Julian Brown wrote: > In C/C++, it's not permitted to mix full struct-type accesses and > "attach/detach" accesses to the same struct's members within a single > directive. This patch adds the same restriction to the Fortran front-end, > and adjusts some existing tests to match. A test is also added for > detection of duplicate derived-type member accesses within a directive. > > Tested with offloading to NVPTX. OK for mainline? Looks good do me, with the comments below addressed. > - the old symbol. gfc_new is used in symbol.c to flag new symbols. */ > + the old symbol. gfc_new is used in symbol.c to flag new symbols. > + comp_mark is used to indicate variables which have component accesses > + in OpenACC directive clauses. */ "OpenMP/OpenACC" – assuming … > + /* Allow multiple components of the same (e.g. derived-type) > + variable here. Duplicate components are detected elsewhere. */ > if (openacc && n->expr && n->expr->expr_type == EXPR_VARIABLE) … > + if (ref->type == REF_COMPONENT) > + component_ref_p = true; > + if (openacc && ((!component_ref_p && n->sym->comp_mark) > + || (component_ref_p && n->sym->mark))) > + gfc_error ("Symbol %qs has mixed component and non-component " > + "accesses at %L", n->sym->name, &n->where); To match C/C++, I think it makes sense to apply the same check also for OpenMP. Currently, that's only a nop as the OpenMP 4.5 feature is not yet implemented (and, hence, component access is rejected during parsing) but the condition also applies to OpenMP. Thanks, Tobias
Hi Julian,
please hold back the commit.
Actually, it does not seem to work if one modifies the test
case a bit. The following code compiles – but I think it
should not:
type one
integer i, j
end type
type two
type(one) A, B
end type
type(two) x
!$acc enter data copyin(x%A, x%A%i)
!$acc enter data copyin(x%A, x%A%i, x%A%i)
end
At least the last line has x%A%i twice but it is accepted.
I am not sure whether the line before is valid or not
(x%A%i is copied twice; might make a difference if it were
a pointer but it isn't a pointer here.) – In C/C++, with OpenMP
and, hence, OpenACC, the first one is accepted but the second
one is rejected. C example:
struct one {
int i, j;
};
struct two {
struct one A, B;
};
void foo() {
struct two x;
/* Accepted: x.A + x.A.i, even though both map 'x.A.i'. */
#pragma omp target enter data map(to:x.A, x.A.i)
#pragma acc enter data copyin(x.A, x.A.i)
/* Rejected: 'x.A.i' mapped twice */
#pragma omp target enter data map(to:x.A, x.A.i, x.A.i)
#pragma acc enter data copyin(x.A, x.A.i, x.A.i)
}
Tobias
On Tue, 7 Jan 2020 11:21:38 +0100 Tobias Burnus <tobias@codesourcery.com> wrote: > Hi Julian, > > please hold back the commit. > > Actually, it does not seem to work if one modifies the test > case a bit. The following code compiles – but I think it > should not: > > type one > integer i, j > end type > type two > type(one) A, B > end type > > type(two) x > !$acc enter data copyin(x%A, x%A%i) > !$acc enter data copyin(x%A, x%A%i, x%A%i) > end > > At least the last line has x%A%i twice but it is accepted. This bit is fixed by: https://gcc.gnu.org/ml/gcc-patches/2020-01/msg00512.html OK again if that one goes in? Thanks, Julian
Hi Julian, On 1/10/20 2:51 AM, Julian Brown wrote: >> Actually, it does not seem to work if one modifies the test >> case a bit. The following code compiles – but I think it >> should not: […] >> At least the last line has x%A%i twice but it is accepted. > This bit is fixed by: > https://gcc.gnu.org/ml/gcc-patches/2020-01/msg00512.html (Side remark: that's checked in gimplify_scan_omp_clauses, which is the "Duplicate components are detected elsewhere".) [ https://gcc.gnu.org/ml/gcc-patches/2020-01/msg00088.html ] > OK again if that one goes in? OK from my side – but I think those patches are orthogonal, i.e. this one can go in without depending on the other one(s). Tobias
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index df07eeed8e3..1de5f0d171e 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1592,9 +1592,11 @@ typedef struct gfc_symbol current statement have the mark member nonzero. Of these symbols, symbols with old_symbol equal to NULL are symbols created within the current statement. Otherwise, old_symbol points to a copy of - the old symbol. gfc_new is used in symbol.c to flag new symbols. */ + the old symbol. gfc_new is used in symbol.c to flag new symbols. + comp_mark is used to indicate variables which have component accesses + in OpenACC directive clauses. */ struct gfc_symbol *old_symbol; - unsigned mark:1, gfc_new:1; + unsigned mark:1, comp_mark:1, gfc_new:1; /* The tlink field is used in the front end to carry the module declaration of separate module procedures so that the characteristics diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index c105657a91b..78351b190f7 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -4245,6 +4245,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, for (n = omp_clauses->lists[list]; n; n = n->next) { n->sym->mark = 0; + n->sym->comp_mark = 0; if (n->sym->attr.flavor == FL_VARIABLE || n->sym->attr.proc_pointer || (!code && (!n->sym->attr.dummy || n->sym->ns != ns))) @@ -4310,23 +4311,25 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, && (list != OMP_LIST_REDUCTION || !openacc)) for (n = omp_clauses->lists[list]; n; n = n->next) { - bool array_only_p = true; - /* Disallow duplicate bare variable references and multiple - subarrays of the same array here, but allow multiple components of - the same (e.g. derived-type) variable. For the latter, duplicate - components are detected elsewhere. */ + bool component_ref_p = false; + + /* Allow multiple components of the same (e.g. derived-type) + variable here. Duplicate components are detected elsewhere. */ if (openacc && n->expr && n->expr->expr_type == EXPR_VARIABLE) for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next) - if (ref->type != REF_ARRAY) - { - array_only_p = false; - break; - } - if (array_only_p) + if (ref->type == REF_COMPONENT) + component_ref_p = true; + if (openacc && ((!component_ref_p && n->sym->comp_mark) + || (component_ref_p && n->sym->mark))) + gfc_error ("Symbol %qs has mixed component and non-component " + "accesses at %L", n->sym->name, &n->where); + else if (n->sym->mark) + gfc_error ("Symbol %qs present on multiple clauses at %L", + n->sym->name, &n->where); + else { - if (n->sym->mark) - gfc_error ("Symbol %qs present on multiple clauses at %L", - n->sym->name, &n->where); + if (component_ref_p) + n->sym->comp_mark = 1; else n->sym->mark = 1; } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-2.f90 b/gcc/testsuite/gfortran.dg/goacc/deep-copy-2.f90 similarity index 63% rename from libgomp/testsuite/libgomp.oacc-fortran/deep-copy-2.f90 rename to gcc/testsuite/gfortran.dg/goacc/deep-copy-2.f90 index 35936617b87..1e556b16b47 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-2.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/deep-copy-2.f90 @@ -1,4 +1,4 @@ -! { dg-do run } +! { dg-do compile } ! Test of attach/detach with "acc data", two clauses at once. @@ -14,7 +14,9 @@ program dtype allocate(var%a(1:n)) -!$acc data copy(var) copy(var%a) +!$acc data copy(var) copy(var%a) ! { dg-error "Symbol .var. has mixed component and non-component accesses" } + +!$acc data copy(var%a) copy(var) ! { dg-error "Symbol .var. has mixed component and non-component accesses" } !$acc parallel loop do i = 1,n @@ -22,6 +24,8 @@ program dtype end do !$acc end parallel loop +!$acc end data + !$acc end data do i = 1,n diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-1.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-1.f90 new file mode 100644 index 00000000000..c1bfe586842 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-1.f90 @@ -0,0 +1,15 @@ +! { dg-do compile } + +subroutine foo + type t + integer :: i, j + end type t + + type(t) x + + ! We should reject the duplicate reference here. +!$acc enter data copyin(x%i, x%i) +! { dg-error ".x.i. appears more than once in map clauses" "" { target "*-*-*" } 11 } + + +end diff --git a/libgomp/testsuite/libgomp.oacc-fortran/classtypes-1.f95 b/libgomp/testsuite/libgomp.oacc-fortran/classtypes-1.f95 index f16f42fc3af..c5f0ffffb02 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/classtypes-1.f95 +++ b/libgomp/testsuite/libgomp.oacc-fortran/classtypes-1.f95 @@ -31,7 +31,8 @@ program main myvar%p%p(i) = -1.0 end do -!$acc enter data copyin(myvar, myvar%p) create(myvar%p%p) +!$acc enter data copyin(myvar) +!$acc enter data copyin(myvar%p) create(myvar%p%p) !$acc parallel loop present(myvar%p%p) do i=1,100 @@ -39,7 +40,8 @@ program main end do !$acc end parallel loop -!$acc exit data copyout(myvar%p%p) delete(myvar, myvar%p) +!$acc exit data copyout(myvar%p%p) delete(myvar%p) +!$acc exit data delete(myvar) do i=1,100 if (myvar%p%p(i) .ne. i * 2) stop 1 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90 index 667d944fecb..edb6b8d61a4 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90 @@ -16,12 +16,14 @@ program dtype allocate(var%a(1:n)) allocate(var%b(1:n)) -!$acc parallel loop copy(var) copy(var%a(1:n)) copy(var%b(1:n)) +!$acc data copy(var) +!$acc parallel loop copy(var%a(1:n)) copy(var%b(1:n)) do i = 1,n var%a(i) = i var%b(i) = i end do !$acc end parallel loop +!$acc end data do i = 1,n if (i .ne. var%a(i)) stop 1
Hi, This patch arose from discussion of the manual deep copy patch here: https://gcc.gnu.org/ml/gcc-patches/2019-12/msg01439.html In C/C++, it's not permitted to mix full struct-type accesses and "attach/detach" accesses to the same struct's members within a single directive. This patch adds the same restriction to the Fortran front-end, and adjusts some existing tests to match. A test is also added for detection of duplicate derived-type member accesses within a directive. Tested with offloading to NVPTX. OK for mainline? Thanks, Julian ChangeLog 2020-01-04 Julian Brown <julian@codesourcery.com> gcc/fortran/ * gfortran.h (gfc_symbol): Add comp_mark bitfield. * openmp.c (resolve_omp_clauses): Disallow mixed component and full-derived-type accesses to the same variable within a single directive. libgomp/ * testsuite/libgomp.oacc-fortran/deep-copy-2.f90: Remove test from here. * testsuite/libgomp.oacc-fortran/deep-copy-3.f90: Don't use mixed component/non-component variable refs in a single directive. * testsuite/libgomp.oacc-fortran/classtypes-1.f95: Likewise. gcc/testsuite/ * gfortran.dg/goacc/deep-copy-2.f90: Add test here. Make a compilation test, and expect rejection of mixed component/non-component accesses. * gfortran.dg/goacc/mapping-tests-1.f90: New test. --- gcc/fortran/gfortran.h | 6 ++-- gcc/fortran/openmp.c | 31 ++++++++++--------- .../gfortran.dg/goacc}/deep-copy-2.f90 | 8 +++-- .../gfortran.dg/goacc/mapping-tests-1.f90 | 15 +++++++++ .../libgomp.oacc-fortran/classtypes-1.f95 | 6 ++-- .../libgomp.oacc-fortran/deep-copy-3.f90 | 4 ++- 6 files changed, 49 insertions(+), 21 deletions(-) rename {libgomp/testsuite/libgomp.oacc-fortran => gcc/testsuite/gfortran.dg/goacc}/deep-copy-2.f90 (63%) create mode 100644 gcc/testsuite/gfortran.dg/goacc/mapping-tests-1.f90