@@ -4313,7 +4313,8 @@ c_omp_address_inspector::expand_array_base (tree c,
/* The code handling "firstprivatize_array_bases" in gimplify.cc is
relevant here. What do we need to create for arrays at this
stage? (This condition doesn't feel quite right. FIXME?) */
- if (!target
+ if (openmp
+ && !target
&& (TREE_CODE (TREE_TYPE (addr_tokens[i + 1]->expr))
== ARRAY_TYPE))
break;
@@ -4324,7 +4325,7 @@ c_omp_address_inspector::expand_array_base (tree c,
virtual_origin);
tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr);
c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
- if (decl_p && target)
+ if (decl_p && (!openmp || target))
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
else
{
@@ -4375,9 +4376,11 @@ c_omp_address_inspector::expand_array_base (tree c,
tree data_addr = omp_accessed_addr (addr_tokens, last_access, expr);
c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
/* For OpenACC, use FIRSTPRIVATE_POINTER for decls even on non-compute
- regions (e.g. "acc data" constructs). It'll be removed anyway in
- gimplify.cc, but doing it this way maintains diagnostic
- behaviour. */
+ regions (e.g. "acc data" constructs). It is used during "lexical
+ inheritance" of mapping clauses on enclosed target
+ (parallel/serial/kernels) regions, i.e. creating "present" mappings
+ for sections of pointer-based arrays. It's also used for
+ diagnostics. */
if (decl_p && (target || !openmp) && !chain_p && !declare_target_p)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
else
@@ -218,6 +218,48 @@ enum gimplify_defaultmap_kind
GDMK_POINTER
};
+/* Used for topological sorting of mapping groups. UNVISITED means we haven't
+ started processing the group yet. The TEMPORARY mark is used when we first
+ encounter a group on a depth-first traversal, and the PERMANENT mark is used
+ when we have processed all the group's children (i.e. all the base pointers
+ referred to by the group's mapping nodes, recursively). */
+
+enum omp_tsort_mark {
+ UNVISITED,
+ TEMPORARY,
+ PERMANENT
+};
+
+/* A group of OMP_CLAUSE_MAP nodes that correspond to a single "map"
+ clause. */
+
+struct omp_mapping_group {
+ tree *grp_start;
+ tree grp_end;
+ omp_tsort_mark mark;
+ /* If we've removed the group but need to reindex, mark the group as
+ deleted. */
+ bool deleted;
+ /* The group points to an already-created "GOMP_MAP_STRUCT
+ GOMP_MAP_ATTACH_DETACH" pair. */
+ bool reprocess_struct;
+ /* The group should use "zero-length" allocations for pointers that are not
+ mapped "to" on the same directive. */
+ bool fragile;
+ struct omp_mapping_group *sibling;
+ struct omp_mapping_group *next;
+
+ omp_mapping_group (tree *_start, tree _end)
+ : grp_start (_start), grp_end (_end), mark (UNVISITED), deleted (false),
+ reprocess_struct (false), fragile (false), sibling (NULL), next (NULL)
+ {
+ }
+
+ omp_mapping_group ()
+ {
+ }
+};
+
struct gimplify_omp_ctx
{
struct gimplify_omp_ctx *outer_context;
@@ -239,6 +281,7 @@ struct gimplify_omp_ctx
bool in_for_exprs;
bool ompacc;
int defaultmap[5];
+ hash_map<tree, omp_mapping_group *> *decl_data_clause;
};
struct privatize_reduction
@@ -473,6 +516,7 @@ new_omp_context (enum omp_region_type region_type)
c->defaultmap[GDMK_AGGREGATE] = GOVD_MAP;
c->defaultmap[GDMK_ALLOCATABLE] = GOVD_MAP;
c->defaultmap[GDMK_POINTER] = GOVD_MAP;
+ c->decl_data_clause = NULL;
return c;
}
@@ -485,6 +529,7 @@ delete_omp_context (struct gimplify_omp_ctx *c)
splay_tree_delete (c->variables);
delete c->privatized_types;
c->loop_iter_var.release ();
+ delete c->decl_data_clause;
XDELETE (c);
}
@@ -8988,18 +9033,6 @@ extract_base_bit_offset (tree base, poly_int64 *bitposp,
return base;
}
-/* Used for topological sorting of mapping groups. UNVISITED means we haven't
- started processing the group yet. The TEMPORARY mark is used when we first
- encounter a group on a depth-first traversal, and the PERMANENT mark is used
- when we have processed all the group's children (i.e. all the base pointers
- referred to by the group's mapping nodes, recursively). */
-
-enum omp_tsort_mark {
- UNVISITED,
- TEMPORARY,
- PERMANENT
-};
-
/* Hash for trees based on operand_equal_p. Like tree_operand_hash
but ignores side effects in the equality comparisons. */
@@ -9016,26 +9049,6 @@ tree_operand_hash_no_se::equal (const value_type &t1,
return operand_equal_p (t1, t2, OEP_MATCH_SIDE_EFFECTS);
}
-/* A group of OMP_CLAUSE_MAP nodes that correspond to a single "map"
- clause. */
-
-struct omp_mapping_group {
- tree *grp_start;
- tree grp_end;
- omp_tsort_mark mark;
- /* If we've removed the group but need to reindex, mark the group as
- deleted. */
- bool deleted;
- /* The group points to an already-created "GOMP_MAP_STRUCT
- GOMP_MAP_ATTACH_DETACH" pair. */
- bool reprocess_struct;
- /* The group should use "zero-length" allocations for pointers that are not
- mapped "to" on the same directive. */
- bool fragile;
- struct omp_mapping_group *sibling;
- struct omp_mapping_group *next;
-};
-
DEBUG_FUNCTION void
debug_mapping_group (omp_mapping_group *grp)
{
@@ -9276,16 +9289,7 @@ omp_gather_mapping_groups_1 (tree *list_p, vec<omp_mapping_group> *groups,
continue;
tree *grp_last_p = omp_group_last (cp);
- omp_mapping_group grp;
-
- grp.grp_start = cp;
- grp.grp_end = *grp_last_p;
- grp.mark = UNVISITED;
- grp.sibling = NULL;
- grp.deleted = false;
- grp.reprocess_struct = false;
- grp.fragile = false;
- grp.next = NULL;
+ omp_mapping_group grp (cp, *grp_last_p);
groups->safe_push (grp);
cp = grp_last_p;
@@ -12267,6 +12271,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
break;
}
+ if (code == OACC_DATA && *grp_start_p != grp_end)
+ {
+ if (!ctx->decl_data_clause)
+ ctx->decl_data_clause = new hash_map<tree, omp_mapping_group *>;
+
+ omp_mapping_group *grp
+ = new omp_mapping_group (grp_start_p, grp_end);
+
+ gcc_assert (DECL_P (decl));
+
+ ctx->decl_data_clause->put (decl, grp);
+ }
flags = GOVD_MAP | GOVD_EXPLICIT;
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
@@ -12953,11 +12969,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
gcc_unreachable ();
}
- if (code == OACC_DATA
- && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
- remove = true;
if (remove)
*list_p = OMP_CLAUSE_CHAIN (c);
else
@@ -13098,6 +13109,52 @@ struct gimplify_adjust_omp_clauses_data
gimple_seq *pre_p;
};
+/* For OpenACC offload regions, the implicit data mappings for arrays must
+ respect explicit data clauses set by a containing acc data region.
+ Specifically, an array section on the data clause must be transformed into
+ an equivalent PRESENT mapping on the inner offload region.
+ This function returns a pointer to a mapping group if an array slice of DECL
+ is specified on a lexically-enclosing data construct, or returns NULL
+ otherwise. */
+
+static omp_mapping_group *
+gomp_oacc_needs_data_present (tree decl)
+{
+ gimplify_omp_ctx *ctx = NULL;
+
+ if (gimplify_omp_ctxp->region_type != ORT_ACC_PARALLEL
+ && gimplify_omp_ctxp->region_type != ORT_ACC_SERIAL
+ && gimplify_omp_ctxp->region_type != ORT_ACC_KERNELS)
+ return NULL;
+
+ if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE
+ && TREE_CODE (TREE_TYPE (decl)) != POINTER_TYPE
+ && TREE_CODE (TREE_TYPE (decl)) != RECORD_TYPE
+ && (TREE_CODE (TREE_TYPE (decl)) != POINTER_TYPE
+ || TREE_CODE (TREE_TYPE (TREE_TYPE (decl))) != ARRAY_TYPE))
+ return NULL;
+
+ decl = get_base_address (decl);
+
+ for (ctx = gimplify_omp_ctxp->outer_context; ctx; ctx = ctx->outer_context)
+ {
+ splay_tree_node on
+ = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
+
+ if (ctx->region_type == ORT_ACC_DATA
+ && on
+ && (((int) on->value) & GOVD_EXPLICIT)
+ && ctx->decl_data_clause != NULL)
+ {
+ omp_mapping_group **pgrp = ctx->decl_data_clause->get (decl);
+ if (pgrp)
+ return *pgrp;
+ }
+ }
+
+ return NULL;
+}
+
/* For all variables that were not actually used within the context,
remove PRIVATE, SHARED, and FIRSTPRIVATE clauses. */
@@ -13219,6 +13276,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
clause = build_omp_clause (input_location, code);
OMP_CLAUSE_DECL (clause) = decl;
OMP_CLAUSE_CHAIN (clause) = chain;
+ omp_mapping_group *outer_grp;
if (private_debug)
OMP_CLAUSE_PRIVATE_DEBUG (clause) = 1;
else if (code == OMP_CLAUSE_PRIVATE && (flags & GOVD_PRIVATE_OUTER_REF))
@@ -13227,6 +13285,58 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
&& (flags & GOVD_WRITTEN) == 0
&& omp_shared_to_firstprivate_optimizable_decl_p (decl))
OMP_CLAUSE_SHARED_READONLY (clause) = 1;
+ else if ((gimplify_omp_ctxp->region_type & ORT_ACC) != 0
+ && (code == OMP_CLAUSE_MAP || code == OMP_CLAUSE_FIRSTPRIVATE)
+ && (outer_grp = gomp_oacc_needs_data_present (decl)))
+ {
+ if (code == OMP_CLAUSE_FIRSTPRIVATE)
+ /* Oops, we have the wrong type of clause. Rebuild it. */
+ clause = build_omp_clause (OMP_CLAUSE_LOCATION (clause),
+ OMP_CLAUSE_MAP);
+
+ tree mapping = *outer_grp->grp_start;
+
+ OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_PRESENT);
+ OMP_CLAUSE_DECL (clause) = unshare_expr (OMP_CLAUSE_DECL (mapping));
+ OMP_CLAUSE_SIZE (clause) = unshare_expr (OMP_CLAUSE_SIZE (mapping));
+
+ /* Copy subsequent nodes (that are part of the mapping group) after the
+ initial one from the outer "acc data" directive -- "pointer" nodes,
+ including firstprivate_reference, pointer sets, etc. */
+
+ tree ptr = OMP_CLAUSE_CHAIN (mapping);
+ tree *ins = &OMP_CLAUSE_CHAIN (clause);
+ tree sentinel = OMP_CLAUSE_CHAIN (outer_grp->grp_end);
+ for (; ptr && ptr != sentinel; ptr = OMP_CLAUSE_CHAIN (ptr))
+ {
+ tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (clause),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (nc, OMP_CLAUSE_MAP_KIND (ptr));
+ OMP_CLAUSE_DECL (nc) = unshare_expr (OMP_CLAUSE_DECL (ptr));
+ OMP_CLAUSE_SIZE (nc) = unshare_expr (OMP_CLAUSE_SIZE (ptr));
+ *ins = nc;
+ ins = &OMP_CLAUSE_CHAIN (nc);
+ }
+
+ *ins = chain;
+
+ gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
+ gimplify_omp_ctxp = ctx->outer_context;
+ for (ptr = clause; ptr != chain; ptr = OMP_CLAUSE_CHAIN (ptr))
+ {
+ /* The condition is specifically to not gimplify here if we have a
+ DECL_P with a DECL_VALUE_EXPR -- i.e. a VLA, or variable-sized
+ array section. If we do, omp-low.cc does not see the DECL_P it
+ expects here for e.g. firstprivate_pointer or
+ firstprivate_reference. */
+ if (!DECL_P (OMP_CLAUSE_DECL (ptr)))
+ gimplify_expr (&OMP_CLAUSE_DECL (ptr), pre_p, NULL,
+ is_gimple_lvalue, fb_lvalue);
+ gimplify_expr (&OMP_CLAUSE_SIZE (ptr), pre_p, NULL,
+ is_gimple_val, fb_rvalue);
+ }
+ gimplify_omp_ctxp = ctx;
+ }
else if (code == OMP_CLAUSE_FIRSTPRIVATE && (flags & GOVD_EXPLICIT) == 0)
OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (clause) = 1;
else if (code == OMP_CLAUSE_MAP && (flags & GOVD_MAP_0LEN_ARRAY) != 0)
@@ -13689,16 +13799,12 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
case OMP_TARGET:
break;
case OACC_DATA:
- if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE)
- break;
- goto check_firstprivate;
case OACC_ENTER_DATA:
case OACC_EXIT_DATA:
case OMP_TARGET_DATA:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
case OACC_HOST_DATA:
- check_firstprivate:
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_FIRSTPRIVATE_REFERENCE))
@@ -21,6 +21,4 @@ int main(int argc, char *argv[])
}
// { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:b\\\[0\\\] \\\[len: 400\\\]\\) map\\(to:a\\\[0\\\] \\\[len: 400\\\]\\)" 1 "gimple" } }
-/* This isn't expected to work while the "lexical inheritance" support is
- reverted. */
-// { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(force_present:b\\\[0\\\] \\\[len: 400\\\]\\) map.alloc:b \\\[pointer assign, bias: 0\\\]\\) map\\(force_present:a\\\[0\\\] \\\[len: 400\\\]\\) map\\(alloc:a \\\[pointer assign, bias: 0\\\]\\)" 0 "gimple" } }
+// { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(force_present:b\\\[0\\\] \\\[len: 400\\\]\\) map\\(firstprivate:b \\\[pointer assign, bias: 0\\\]\\) map\\(force_present:a\\\[0\\\] \\\[len: 400\\\]\\) map\\(firstprivate:a \\\[pointer assign, bias: 0\\\]\\)" 1 "gimple" } }
new file mode 100644
@@ -0,0 +1,35 @@
+! Test if implicitly determined data clauses work with an
+! assumed-sized array variable. Note that the array variable, 'a',
+! has been explicitly copied in and out via acc enter data and acc
+! exit data, respectively.
+
+! This does not appear to be supported by the OpenACC standard as of version
+! 3.0. Check for an appropriate error message.
+
+program test
+ implicit none
+
+ integer, parameter :: n = 100
+ integer a(n), i
+
+ call dtest (a, n)
+
+ do i = 1, n
+ if (a(i) /= i) call abort
+ end do
+end program test
+
+subroutine dtest (a, n)
+ integer i, n
+ integer a(*)
+
+ !$acc enter data copyin(a(1:n))
+
+ !$acc parallel loop
+! { dg-error {implicit mapping of assumed size array 'a'} "" { target *-*-* } .-1 }
+ do i = 1, n
+ a(i) = i
+ end do
+
+ !$acc exit data copyout(a(1:n))
+end subroutine dtest
@@ -19,5 +19,4 @@ program test
end program test
! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(tofrom:data\\\[\_\[0-9\]+\\\] \\\[len: _\[0-9\]+\\\]\\) map\\(alloc:data \\\[pointer assign, bias: _\[0-9\]+\\\]\\)" 1 "gimple" } }
-! Disable for now
-! { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(force_present:data\\\[D\\.\[0-9\]+\\\] \\\[len: D\\.\[0-9\]+\\\]\\) map\\(alloc:data \\\[pointer assign, bias: D\\.\[0-9\]+\\\]\\)" 0 "gimple" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(force_present:data\\\[D\\.\[0-9\]+\\\] \\\[len: D\\.\[0-9\]+\\\]\\) map\\(alloc:data \\\[pointer assign, bias: D\\.\[0-9\]+\\\]\\)" 1 "gimple" } }
@@ -32,5 +32,3 @@ main (int argc, char* argv[])
return 0;
}
-
-/* { dg-xfail-run-if "PR70828" { ! openacc_host_selected } } */
@@ -25,5 +25,3 @@ main ()
return 0;
}
-
-/* { dg-xfail-run-if "PR70828" { ! openacc_host_selected } } */
@@ -29,5 +29,3 @@ program test
end if
end do
end program test
-
-! { dg-xfail-run-if "PR70828" { ! openacc_host_selected } }
@@ -32,5 +32,3 @@ program test
end if
end do
end program test
-
-! { dg-xfail-run-if "PR70828" { ! openacc_host_selected } }
@@ -29,5 +29,3 @@ program test
end if
end do
end program test
-
-! { dg-xfail-run-if "PR70828" { ! openacc_host_selected } }
@@ -27,5 +27,3 @@ program test
end if
end do
end program test
-
-! { dg-xfail-run-if "PR70828" { ! openacc_host_selected } }
@@ -26,5 +26,3 @@ program test
end if
end do
end program test
-
-! { dg-xfail-run-if "PR70828" { ! openacc_host_selected } }
@@ -22,5 +22,3 @@ program test
end if
end do
end program test
-
-! { dg-xfail-run-if "PR70828" { ! openacc_host_selected } }
This patch reimplements "lexical inheritance" for OpenACC offload regions inside "data" regions, allowing e.g. this to work: int *ptr; [...] #pragma acc data copyin(ptr[10:2]) { #pragma acc parallel { ... } } here, the "copyin" is mirrored on the inner "acc parallel" as "present(ptr[10:2])" -- allowing code within the parallel to use that section of the array even though the mapping is implicit. In terms of implementation, this works by expanding mapping nodes for "acc data" to include pointer mappings that might be needed by inner offload regions. The resulting mapping group is then copied to the inner offload region as needed, rewriting the first node to "force_present". The pointer mapping nodes are then removed from the "acc data" later during gimplification. For OpenMP, pointer mapping nodes on equivalent "omp data" regions are not needed, so remain suppressed during expansion. 2023-06-16 Julian Brown <julian@codesourcery.com> gcc/c-family/ * c-omp.cc (c_omp_address_inspector::expand_array_base): Don't omit pointer nodes for OpenACC. gcc/ * gimplify.cc (omp_tsort_mark, omp_mapping_group): Move before gimplify_omp_ctx. Add constructor to omp_mapping_group. (gimplify_omp_ctx): Add DECL_DATA_CLAUSE field. (new_omp_context, delete_omp_context): Initialise and free above field. (omp_gather_mapping_groups_1): Use constructor for omp_mapping_group. (gimplify_scan_omp_clauses): Record mappings that might be lexically inherited. Don't remove GOMP_MAP_FIRSTPRIVATE_POINTER/GOMP_MAP_FIRSTPRIVATE_REFERENCE yet. (gomp_oacc_needs_data_present): New function. (gimplify_adjust_omp_clauses_1): Implement lexical inheritance behaviour for OpenACC. (gimplify_adjust_omp_clauses): Remove GOMP_MAP_FIRSTPRIVATE_POINTER/GOMP_MAP_FIRSTPRIVATE_REFERENCE here instead, after lexical inheritance is done. gcc/testsuite/ * c-c++-common/goacc/acc-data-chain.c: Re-enable scan test. * gfortran.dg/goacc/pr70828.f90: Likewise. * gfortran.dg/goacc/assumed-size.f90: New test. libgomp/ * testsuite/libgomp.oacc-c-c++-common/pr70828.c: Un-XFAIL. * testsuite/libgomp.oacc-c-c++-common/pr70828-2.c: Un-XFAIL. * testsuite/libgomp.oacc-fortran/pr70828.f90: Un-XFAIL. * testsuite/libgomp.oacc-fortran/pr70828-2.f90: Un-XFAIL. * testsuite/libgomp.oacc-fortran/pr70828-3.f90: Un-XFAIL. * testsuite/libgomp.oacc-fortran/pr70828-4.f90: Un-XFAIL. * testsuite/libgomp.oacc-fortran/pr70828-5.f90: Un-XFAIL. * testsuite/libgomp.oacc-fortran/pr70828-6.f90: Un-XFAIL. --- gcc/c-family/c-omp.cc | 13 +- gcc/gimplify.cc | 208 +++++++++++++----- .../c-c++-common/goacc/acc-data-chain.c | 4 +- .../gfortran.dg/goacc/assumed-size.f90 | 35 +++ gcc/testsuite/gfortran.dg/goacc/pr70828.f90 | 3 +- .../libgomp.oacc-c-c++-common/pr70828-2.c | 2 - .../libgomp.oacc-c-c++-common/pr70828.c | 2 - .../libgomp.oacc-fortran/pr70828-2.f90 | 2 - .../libgomp.oacc-fortran/pr70828-3.f90 | 2 - .../libgomp.oacc-fortran/pr70828-4.f90 | 2 - .../libgomp.oacc-fortran/pr70828-5.f90 | 2 - .../libgomp.oacc-fortran/pr70828-6.f90 | 2 - .../libgomp.oacc-fortran/pr70828.f90 | 2 - 13 files changed, 202 insertions(+), 77 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/goacc/assumed-size.f90