https://gcc.gnu.org/g:ed4a45ac95375fac7166cff38dc9a81bd3868abd
commit ed4a45ac95375fac7166cff38dc9a81bd3868abd Author: Julian Brown <[email protected]> Date: Wed Apr 23 03:09:01 2025 +0000 OpenACC: "declare create" fixes wrt. "allocatable" variables This patch fixes a case revealed by the previous patch where a synthetic "acc data" region created for a "declare create" variable could interact strangely with lexical inheritance behaviour. In fact, it doesn't seem right to create the "acc data" region for allocatable variables at all -- doing so means that a data region is likely to be created for an unallocated variable. The fix is not to add such variables to the synthetic "acc data" region at all, and defer to the code that performs "enter data"/"exit data" for them when allocated/deallocated on the host instead. Then, "declare create" variables are implicitly turned into "present" clauses on in-scope offload regions. gcc/fortran/ * trans-openmp.cc (gfc_omp_finish_clause): Handle "declare create" for scalar allocatable variables. (gfc_trans_omp_clauses): Don't include allocatable vars in synthetic "acc data" region created for "declare create" variables. Mark such variables with the "oacc declare create" attribute instead. Don't create ALWAYS_POINTER mapping for target-to-host updates of declare create variables. (gfc_trans_oacc_declare): Handle empty clause list. gcc/ * gimplify.cc (gimplify_adjust_omp_clauses_1): Handle "oacc declare create" attribute. gcc/testsuite/ * c-c++-common/goacc/readonly-1.c: Adjust patterns. libgomp/ * testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90: Remove xfails. * testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90: Remove xfails. * testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90: Remove xfails. * testsuite/libgomp.oacc-fortran/declare-create-1.f90: New test. * testsuite/libgomp.oacc-fortran/declare-create-2.f90: New test. * testsuite/libgomp.oacc-fortran/declare-create-3.f90: New test. Co-Authored-By: Paul-Antoine Arras <[email protected]> Co-Authored-By: Sandra Loosemore <[email protected]> Diff: --- gcc/fortran/trans-openmp.cc | 45 +++++++++++++++++++--- gcc/gimplify.cc | 8 ++++ gcc/testsuite/c-c++-common/goacc/readonly-1.c | 8 ++-- .../declare-allocatable-1-directive.f90 | 6 +-- .../declare-allocatable-1-runtime.f90 | 6 +-- .../libgomp.oacc-fortran/declare-allocatable-1.f90 | 6 +-- .../libgomp.oacc-fortran/declare-create-1.f90 | 22 +++++++++++ .../libgomp.oacc-fortran/declare-create-2.f90 | 26 +++++++++++++ .../libgomp.oacc-fortran/declare-create-3.f90 | 26 +++++++++++++ 9 files changed, 135 insertions(+), 18 deletions(-) diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 36a3bcd8e14c..c28d937295a6 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -1627,7 +1627,16 @@ gfc_omp_finish_clause (tree c, gimple_seq *pre_p, bool openacc) orig_decl = decl; c4 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (c4, GOMP_MAP_POINTER); + if (openacc + && GFC_DECL_GET_SCALAR_ALLOCATABLE (decl) + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_PRESENT) + /* This allows "declare create" to work for scalar allocatables. The + resulting mapping nodes are: + force_present(*var) firstprivate_pointer(var) + which is the same as an explicit "present" clause gives. */ + OMP_CLAUSE_SET_MAP_KIND (c4, GOMP_MAP_FIRSTPRIVATE_POINTER); + else + OMP_CLAUSE_SET_MAP_KIND (c4, GOMP_MAP_POINTER); OMP_CLAUSE_DECL (c4) = decl; OMP_CLAUSE_SIZE (c4) = size_int (0); decl = build_fold_indirect_ref (decl); @@ -4025,6 +4034,29 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, if (!n->sym->attr.referenced) continue; + /* We do not want to include allocatable vars in a synthetic + "acc data" region created for "!$acc declare create" vars. + Such variables are handled by augmenting allocate/deallocate + statements elsewhere (with + "acc enter data declare_allocate(...)", etc.). */ + if (op == EXEC_OACC_DECLARE + && n->u.map.op == OMP_MAP_ALLOC + && n->sym->attr.allocatable + && n->sym->attr.oacc_declare_create) + { + tree tree_var = gfc_get_symbol_decl (n->sym); + if (!lookup_attribute ("oacc declare create", + DECL_ATTRIBUTES (tree_var))) + DECL_ATTRIBUTES (tree_var) + = tree_cons (get_identifier ("oacc declare create"), + NULL_TREE, DECL_ATTRIBUTES (tree_var)); + /* We might need to turn what would normally be a + "firstprivate" mapping into a "present" mapping. For the + latter, we need the decl to be addressable. */ + TREE_ADDRESSABLE (tree_var) = 1; + continue; + } + location_t map_loc = gfc_get_location (&n->where); bool always_modifier = false; tree node = build_omp_clause (map_loc, OMP_CLAUSE_MAP); @@ -4255,7 +4287,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, else if (op == EXEC_OMP_TARGET_EXIT_DATA) gmk = GOMP_MAP_RELEASE; else if (GFC_DECL_GET_SCALAR_ALLOCATABLE (decl) - && n->sym->attr.oacc_declare_create) + && n->sym->attr.oacc_declare_create + && n->u.map.op != OMP_MAP_FORCE_FROM) { if (clauses->update_allocatable) gmk = GOMP_MAP_ALWAYS_POINTER; @@ -9363,10 +9396,12 @@ gfc_trans_oacc_declare (gfc_code *code) gfc_start_block (&block); oacc_clauses = gfc_trans_omp_clauses (&block, code->ext.oacc_declare->clauses, - code->loc, false, true); + code->loc, false, true, + EXEC_OACC_DECLARE); stmt = gfc_trans_omp_code (code->block->next, true); - stmt = build2_loc (input_location, construct_code, void_type_node, stmt, - oacc_clauses); + if (oacc_clauses) + stmt = build2_loc (input_location, construct_code, void_type_node, stmt, + oacc_clauses); gfc_add_expr_to_block (&block, stmt); return gfc_finish_block (&block); diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 31e4b4d97b0c..00fb34f71827 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -14647,6 +14647,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) g->have_offload = true; } } + if (lookup_attribute ("oacc declare create", DECL_ATTRIBUTES (decl))) + flags |= GOVD_MAP_FORCE_PRESENT; } else if (flags & GOVD_SHARED) { @@ -14686,6 +14688,12 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) "%<target%> construct", decl); return 0; } + if (lookup_attribute ("oacc declare create", DECL_ATTRIBUTES (decl))) + { + code = OMP_CLAUSE_MAP; + flags &= ~GOVD_FIRSTPRIVATE; + flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT; + } } else if (flags & GOVD_LASTPRIVATE) code = OMP_CLAUSE_LASTPRIVATE; diff --git a/gcc/testsuite/c-c++-common/goacc/readonly-1.c b/gcc/testsuite/c-c++-common/goacc/readonly-1.c index 300464c92e3f..bd4d2294c2a5 100644 --- a/gcc/testsuite/c-c++-common/goacc/readonly-1.c +++ b/gcc/testsuite/c-c++-common/goacc/readonly-1.c @@ -51,14 +51,14 @@ int main (void) /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */ /* { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */ /* { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */ -/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */ -/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */ +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */ +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */ /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */ /* { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */ /* { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */ -/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */ -/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */ +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */ +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */ /* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\);$" 4 "original" } } */ /* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\);$" 4 "original" } } */ diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90 index bdeabca3eb50..6e53dc563b24 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90 @@ -212,9 +212,9 @@ program test !$acc exit data delete (b) deallocate (b) end program test ! { dg-line l[incr c] } -! { dg-bogus {note: variable 'overflow\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO n/a} { xfail *-*-* } l$c } -! { dg-bogus {note: variable 'not_prev_allocated\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO n/a} { xfail *-*-* } l$c } -! { dg-bogus {note: variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} {TODO n/a} { xfail *-*-* } l$c } +! { dg-bogus {note: variable 'overflow\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } +! { dg-bogus {note: variable 'not_prev_allocated\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } +! { dg-bogus {note: variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} {} { target *-*-* } l$c } ! Set each element in array 'b' at index i to i*2. diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90 index a3e17c82828d..0072827d34c0 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90 @@ -212,9 +212,9 @@ program test call acc_delete (b) deallocate (b) end program test ! { dg-line l[incr c] } -! { dg-bogus {note: variable 'overflow\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO n/a} { xfail *-*-* } l$c } -! { dg-bogus {note: variable 'not_prev_allocated\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO n/a} { xfail *-*-* } l$c } -! { dg-bogus {note: variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} {TODO n/a} { xfail *-*-* } l$c } +! { dg-bogus {note: variable 'overflow\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } +! { dg-bogus {note: variable 'not_prev_allocated\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } +! { dg-bogus {note: variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} {} { target *-*-* } l$c } ! Set each element in array 'b' at index i to i*2. diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90 index 7220661b54aa..ab6ff753c63e 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90 @@ -202,9 +202,9 @@ program test deallocate (b) end program test ! { dg-line l[incr c] } -! { dg-bogus {note: variable 'overflow\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO n/a} { xfail *-*-* } l$c } -! { dg-bogus {note: variable 'not_prev_allocated\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO n/a} { xfail *-*-* } l$c } -! { dg-bogus {note: variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} {TODO n/a} { xfail *-*-* } l$c } +! { dg-bogus {note: variable 'overflow\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } +! { dg-bogus {note: variable 'not_prev_allocated\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } +! { dg-bogus {note: variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} {} { target *-*-* } l$c } ! Set each element in array 'b' at index i to i*2. diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-create-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-create-1.f90 new file mode 100644 index 000000000000..057b5eb958ad --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-create-1.f90 @@ -0,0 +1,22 @@ +! { dg-do run } + +module m +integer :: mint +!$acc declare create (mint) +end module m + +program p +use m + +mint = 0 + +!$acc serial +! { dg-warning {using .vector_length \(32\)., ignoring 1} "" { target openacc_nvidia_accel_selected } .-1 } +mint = 5 +!$acc end serial + +!$acc update host(mint) + +if (mint.ne.5) stop 1 + +end program p diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-create-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-create-2.f90 new file mode 100644 index 000000000000..dd7c9798fbae --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-create-2.f90 @@ -0,0 +1,26 @@ +! { dg-do run } + +module m +integer, allocatable :: mint +!$acc declare create (mint) +end module m + +program p +use m + +allocate(mint) + +mint = 0 + +!$acc serial +! { dg-warning {using .vector_length \(32\)., ignoring 1} "" { target openacc_nvidia_accel_selected } .-1 } +mint = 5 +!$acc end serial + +!$acc update host(mint) + +if (mint.ne.5) stop 1 + +deallocate(mint) + +end program p diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-create-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-create-3.f90 new file mode 100644 index 000000000000..7cceaa5f8a3f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-create-3.f90 @@ -0,0 +1,26 @@ +! { dg-do run } + +module m +integer, allocatable :: mint(:) +!$acc declare create (mint) +end module m + +program p +use m + +allocate(mint(1:20)) + +mint = 0 + +!$acc serial +! { dg-warning {using .vector_length \(32\)., ignoring 1} "" { target openacc_nvidia_accel_selected } .-1 } +mint = 5 +!$acc end serial + +!$acc update host(mint) + +if (any(mint.ne.5)) stop 1 + +deallocate(mint) + +end program p
