| Message ID | 20180828151919.576c636c@squid.athome |
|---|---|
| State | New |
| Headers | show |
| Series | [OpenACC] (1/2) Fix implicit mapping for array slices on lexically-enclosing data constructs (PR70828) | expand |
On Tue, Aug 28, 2018 at 03:19:19PM -0400, Julian Brown wrote: > 2018-08-28 Julian Brown <julian@codesourcery.com> > Cesar Philippidis <cesar@codesourcery.com> > > PR middle-end/70828 > > gcc/ > * gimplify.c (gimplify_omp_ctx): Add decl_data_clause hash map. > (new_omp_context): Initialise above. > (delete_omp_context): Delete above. > (gimplify_scan_omp_clauses): Scan for array mappings on data constructs, > and record in above map. > (gomp_needs_data_present): New function. > (gimplify_adjust_omp_clauses_1): Handle data mappings (e.g. array > slices) declared in lexically-enclosing data constructs. > * omp-low.c (lower_omp_target): Allow decl for bias not to be present > in omp context. > > gcc/testsuite/ > * c-c++-common/goacc/acc-data-chain.c: New test. > * gfortran.dg/goacc/pr70828.f90: New test. > * gfortran.dg/goacc/pr70828-2.f90: New test. > > libgomp/ > * testsuite/libgomp.oacc-c-c++-common/pr70828.c: New test. > * testsuite/libgomp.oacc-fortran/implicit_copy.f90: New test. > * testsuite/libgomp.oacc-fortran/pr70828.f90: New test. > * testsuite/libgomp.oacc-fortran/pr70828-2.f90: New test. > * testsuite/libgomp.oacc-fortran/pr70828-3.f90: New test. > * testsuite/libgomp.oacc-fortran/pr70828-5.f90: New test. > --- a/gcc/gimplify.c > +++ b/gcc/gimplify.c > @@ -191,6 +191,7 @@ struct gimplify_omp_ctx > bool target_map_scalars_firstprivate; > bool target_map_pointers_as_0len_arrays; > bool target_firstprivatize_array_bases; > + hash_map<tree, std::pair<tree, tree> > *decl_data_clause; > }; > > static struct gimplify_ctx *gimplify_ctxp; > @@ -413,6 +414,7 @@ new_omp_context (enum omp_region_type region_type) > c->default_kind = OMP_CLAUSE_DEFAULT_SHARED; > else > c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED; > + c->decl_data_clause = new hash_map<tree, std::pair<tree, tree> >; Not really happy about creating this unconditionally. Can you leave it NULL by default and only initialize for contexts where it will be needed? > @@ -7793,8 +7796,21 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, > case OMP_TARGET: > break; > case OACC_DATA: > - if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE) > - break; > + { > + tree nextc = OMP_CLAUSE_CHAIN (c); > + if (nextc > + && OMP_CLAUSE_CODE (nextc) == OMP_CLAUSE_MAP > + && (OMP_CLAUSE_MAP_KIND (nextc) > + == GOMP_MAP_FIRSTPRIVATE_POINTER > + || OMP_CLAUSE_MAP_KIND (nextc) == GOMP_MAP_POINTER)) > + { > + tree base_addr = OMP_CLAUSE_DECL (nextc); > + ctx->decl_data_clause->put (base_addr, > + std::make_pair (unshare_expr (c), unshare_expr (nextc))); Don't like the wrapping here, can you just split it up: std::pair<tree, tree> p = std::make_pair (unshare_expr (c), unshare_expr (nextc)); ctx->decl_data_clause->put (base_addr, p); or similar? > + > +static std::pair<tree, tree> * > +gomp_needs_data_present (tree decl) Would be helpful to have acc/oacc in the function name. > +{ > + gimplify_omp_ctx *ctx = NULL; > + > + if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE > + && TREE_CODE (TREE_TYPE (decl)) != POINTER_TYPE > + && (TREE_CODE (TREE_TYPE (decl)) != POINTER_TYPE > + || TREE_CODE (TREE_TYPE (TREE_TYPE (decl))) != ARRAY_TYPE)) > + return NULL; > + > + if (gimplify_omp_ctxp->region_type != ORT_ACC_PARALLEL > + && gimplify_omp_ctxp->region_type != ORT_ACC_KERNELS) > + return NULL; And move this test to the top. > --- a/gcc/omp-low.c > +++ b/gcc/omp-low.c > @@ -8411,9 +8411,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) > x = fold_convert_loc (clause_loc, type, x); > if (!integer_zerop (OMP_CLAUSE_SIZE (c))) > { > - tree bias = OMP_CLAUSE_SIZE (c); > - if (DECL_P (bias)) > - bias = lookup_decl (bias, ctx); > + tree bias = OMP_CLAUSE_SIZE (c), remapped_bias; > + if (DECL_P (bias) > + && (remapped_bias = maybe_lookup_decl (bias, ctx))) > + bias = remapped_bias; > bias = fold_convert_loc (clause_loc, sizetype, bias); > bias = fold_build1_loc (clause_loc, NEGATE_EXPR, sizetype, > bias); This is shared with OpenMP and must be conditionalized for OpenACC only. Jakub
On Tue, 4 Dec 2018 15:02:15 +0100 Jakub Jelinek <jakub@redhat.com> wrote: > On Tue, Aug 28, 2018 at 03:19:19PM -0400, Julian Brown wrote: > > 2018-08-28 Julian Brown <julian@codesourcery.com> > > Cesar Philippidis <cesar@codesourcery.com> > > > > PR middle-end/70828 > > > > gcc/ > > * gimplify.c (gimplify_omp_ctx): Add decl_data_clause hash > > map. (new_omp_context): Initialise above. > > (delete_omp_context): Delete above. > > (gimplify_scan_omp_clauses): Scan for array mappings on > > data constructs, and record in above map. > > (gomp_needs_data_present): New function. > > (gimplify_adjust_omp_clauses_1): Handle data mappings (e.g. > > array slices) declared in lexically-enclosing data constructs. > > * omp-low.c (lower_omp_target): Allow decl for bias not to > > be present in omp context. > > > > gcc/testsuite/ > > * c-c++-common/goacc/acc-data-chain.c: New test. > > * gfortran.dg/goacc/pr70828.f90: New test. > > * gfortran.dg/goacc/pr70828-2.f90: New test. > > > > libgomp/ > > * testsuite/libgomp.oacc-c-c++-common/pr70828.c: New test. > > * testsuite/libgomp.oacc-fortran/implicit_copy.f90: New > > test. > > * testsuite/libgomp.oacc-fortran/pr70828.f90: New test. > > * testsuite/libgomp.oacc-fortran/pr70828-2.f90: New test. > > * testsuite/libgomp.oacc-fortran/pr70828-3.f90: New test. > > * testsuite/libgomp.oacc-fortran/pr70828-5.f90: New test. > > > --- a/gcc/gimplify.c > > +++ b/gcc/gimplify.c > > @@ -191,6 +191,7 @@ struct gimplify_omp_ctx > > bool target_map_scalars_firstprivate; > > bool target_map_pointers_as_0len_arrays; > > bool target_firstprivatize_array_bases; > > + hash_map<tree, std::pair<tree, tree> > *decl_data_clause; > > }; > > > > static struct gimplify_ctx *gimplify_ctxp; > > @@ -413,6 +414,7 @@ new_omp_context (enum omp_region_type > > region_type) c->default_kind = OMP_CLAUSE_DEFAULT_SHARED; > > else > > c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED; > > + c->decl_data_clause = new hash_map<tree, std::pair<tree, tree> > > >; > > Not really happy about creating this unconditionally. Can you leave > it NULL by default and only initialize for contexts where it will be > needed? > > > @@ -7793,8 +7796,21 @@ gimplify_scan_omp_clauses (tree *list_p, > > gimple_seq *pre_p, case OMP_TARGET: > > break; > > case OACC_DATA: > > - if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE) > > - break; > > + { > > + tree nextc = OMP_CLAUSE_CHAIN (c); > > + if (nextc > > + && OMP_CLAUSE_CODE (nextc) == OMP_CLAUSE_MAP > > + && (OMP_CLAUSE_MAP_KIND (nextc) > > + == GOMP_MAP_FIRSTPRIVATE_POINTER > > + || OMP_CLAUSE_MAP_KIND (nextc) == > > GOMP_MAP_POINTER)) > > + { > > + tree base_addr = OMP_CLAUSE_DECL (nextc); > > + ctx->decl_data_clause->put (base_addr, > > + std::make_pair (unshare_expr (c), > > unshare_expr (nextc))); > > Don't like the wrapping here, can you just split it up: > std::pair<tree, tree> p > = std::make_pair (unshare_expr (c), > unshare_expr (nextc)); > ctx->decl_data_clause->put (base_addr, p); > or similar? > > > + > > +static std::pair<tree, tree> * > > +gomp_needs_data_present (tree decl) > > Would be helpful to have acc/oacc in the function name. > > +{ > > + gimplify_omp_ctx *ctx = NULL; > > + > > + if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE > > + && TREE_CODE (TREE_TYPE (decl)) != POINTER_TYPE > > + && (TREE_CODE (TREE_TYPE (decl)) != POINTER_TYPE > > + || TREE_CODE (TREE_TYPE (TREE_TYPE (decl))) != > > ARRAY_TYPE)) > > + return NULL; > > + > > + if (gimplify_omp_ctxp->region_type != ORT_ACC_PARALLEL > > + && gimplify_omp_ctxp->region_type != ORT_ACC_KERNELS) > > + return NULL; > > And move this test to the top. > > > --- a/gcc/omp-low.c > > +++ b/gcc/omp-low.c > > @@ -8411,9 +8411,10 @@ lower_omp_target (gimple_stmt_iterator > > *gsi_p, omp_context *ctx) x = fold_convert_loc (clause_loc, type, > > x); if (!integer_zerop (OMP_CLAUSE_SIZE (c))) > > { > > - tree bias = OMP_CLAUSE_SIZE (c); > > - if (DECL_P (bias)) > > - bias = lookup_decl (bias, ctx); > > + tree bias = OMP_CLAUSE_SIZE (c), remapped_bias; > > + if (DECL_P (bias) > > + && (remapped_bias = maybe_lookup_decl > > (bias, ctx))) > > + bias = remapped_bias; > > bias = fold_convert_loc (clause_loc, sizetype, > > bias); bias = fold_build1_loc (clause_loc, NEGATE_EXPR, sizetype, > > bias); > > This is shared with OpenMP and must be conditionalized for OpenACC > only. Thanks for review! How's this version? I took the liberty of fixing the patch for Fortran array-descriptor mappings that use a PSET, also, and adding another test for that functionality. Re-tested with offloading to nvptx. OK? Julian 2018-08-28 Julian Brown <julian@codesourcery.com> Cesar Philippidis <cesar@codesourcery.com> gcc/ * gimplify.c (oacc_array_mapping_info): New struct. (gimplify_omp_ctx): Add decl_data_clause hash map. (new_omp_context): Zero-initialise above. (delete_omp_context): Delete above if allocated. (gimplify_scan_omp_clauses): Scan for array mappings on data constructs, and record in above map. (gomp_oacc_needs_data_present): New function. (gimplify_adjust_omp_clauses_1): Handle data mappings (e.g. array slices) declared in lexically-enclosing data constructs. * omp-low.c (lower_omp_target): Allow decl for bias not to be present in OpenACC context. gcc/testsuite/ * c-c++-common/goacc/acc-data-chain.c: New test. * gfortran.dg/goacc/pr70828.f90: New test. * gfortran.dg/goacc/pr70828-2.f90: New test. libgomp/ * testsuite/libgomp.oacc-c-c++-common/pr70828.c: New test. * testsuite/libgomp.oacc-fortran/implicit_copy.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828-2.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828-3.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828-5.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828-6.f90: New test. Reviewed-by: Jakub Jelinek <jakub@redhat.com> commit 390adf97cfdde951ed1e82fc54d77e34130c70b8 Author: Julian Brown <julian@codesourcery.com> Date: Thu Aug 16 20:02:10 2018 -0700 Inheritance of array sections on data constructs. 2018-08-28 Julian Brown <julian@codesourcery.com> Cesar Philippidis <cesar@codesourcery.com> gcc/ * gimplify.c (oacc_array_mapping_info): New struct. (gimplify_omp_ctx): Add decl_data_clause hash map. (new_omp_context): Zero-initialise above. (delete_omp_context): Delete above if allocated. (gimplify_scan_omp_clauses): Scan for array mappings on data constructs, and record in above map. (gomp_oacc_needs_data_present): New function. (gimplify_adjust_omp_clauses_1): Handle data mappings (e.g. array slices) declared in lexically-enclosing data constructs. * omp-low.c (lower_omp_target): Allow decl for bias not to be present in OpenACC context. gcc/testsuite/ * c-c++-common/goacc/acc-data-chain.c: New test. * gfortran.dg/goacc/pr70828.f90: New test. * gfortran.dg/goacc/pr70828-2.f90: New test. libgomp/ * testsuite/libgomp.oacc-c-c++-common/pr70828.c: New test. * testsuite/libgomp.oacc-fortran/implicit_copy.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828-2.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828-3.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828-5.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828-6.f90: New test. diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 509fc2f..b6a9bfc 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -176,6 +176,17 @@ struct gimplify_ctx unsigned in_switch_expr : 1; }; +/* Used to record clauses representing array slices on data directives that + may affect implicit mapping semantics on enclosed OpenACC parallel/kernels + regions. PSET is used for Fortran array slices with array descriptors, + or NULL otherwise. */ +struct oacc_array_mapping_info +{ + tree mapping; + tree pset; + tree pointer; +}; + struct gimplify_omp_ctx { struct gimplify_omp_ctx *outer_context; @@ -191,6 +202,7 @@ struct gimplify_omp_ctx bool target_map_scalars_firstprivate; bool target_map_pointers_as_0len_arrays; bool target_firstprivatize_array_bases; + hash_map<tree, oacc_array_mapping_info> *decl_data_clause; }; static struct gimplify_ctx *gimplify_ctxp; @@ -413,6 +425,7 @@ new_omp_context (enum omp_region_type region_type) c->default_kind = OMP_CLAUSE_DEFAULT_SHARED; else c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED; + c->decl_data_clause = NULL; return c; } @@ -425,6 +438,8 @@ delete_omp_context (struct gimplify_omp_ctx *c) splay_tree_delete (c->variables); delete c->privatized_types; c->loop_iter_var.release (); + if (c->decl_data_clause) + delete c->decl_data_clause; XDELETE (c); } @@ -7795,8 +7810,41 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_TARGET: break; case OACC_DATA: - if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE) - break; + { + tree base_ptr = OMP_CLAUSE_CHAIN (c); + tree pset = NULL; + if (base_ptr + && OMP_CLAUSE_CODE (base_ptr) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (base_ptr) == GOMP_MAP_TO_PSET) + { + pset = base_ptr; + base_ptr = OMP_CLAUSE_CHAIN (base_ptr); + } + if (base_ptr + && OMP_CLAUSE_CODE (base_ptr) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET + && ((OMP_CLAUSE_MAP_KIND (base_ptr) + == GOMP_MAP_FIRSTPRIVATE_POINTER) + || OMP_CLAUSE_MAP_KIND (base_ptr) == GOMP_MAP_POINTER)) + { + /* If we have an array descriptor, fish the right base + address variable to use out of that (otherwise we'd have + to deconstruct "arr.data" in the subsequent pointer + mapping). */ + tree base_addr = pset ? OMP_CLAUSE_DECL (pset) + : OMP_CLAUSE_DECL (base_ptr); + if (!ctx->decl_data_clause) + ctx->decl_data_clause + = new hash_map<tree, oacc_array_mapping_info>; + oacc_array_mapping_info ai; + ai.mapping = unshare_expr (c); + ai.pset = pset ? unshare_expr (pset) : NULL; + ai.pointer = unshare_expr (base_ptr); + ctx->decl_data_clause->put (base_addr, ai); + } + if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE) + break; + } /* FALLTHRU */ case OMP_TARGET_DATA: case OMP_TARGET_ENTER_DATA: @@ -8695,6 +8743,46 @@ struct gimplify_adjust_omp_clauses_data gimple_seq *pre_p; }; +/* For OpenACC parallel and kernels 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 parallel or + kernels region. This function returns a pointer to an + oacc_array_mapping_info if an array slice of DECL is specified in a + lexically-enclosing data construct, or returns NULL otherwise. */ + +static oacc_array_mapping_info * +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_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) + { + oacc_array_mapping_info *ret; + + if (ctx->region_type != ORT_ACC_DATA) + break; + + if (ctx->decl_data_clause && (ret = ctx->decl_data_clause->get (decl))) + return ret; + } + + return NULL; +} + /* For all variables that were not actually used within the context, remove PRIVATE, SHARED, and FIRSTPRIVATE clauses. */ @@ -8787,6 +8875,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; + oacc_array_mapping_info *array_info; if (private_debug) OMP_CLAUSE_PRIVATE_DEBUG (clause) = 1; else if (code == OMP_CLAUSE_PRIVATE && (flags & GOVD_PRIVATE_OUTER_REF)) @@ -8795,6 +8884,56 @@ 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 ((code == OMP_CLAUSE_MAP || code == OMP_CLAUSE_FIRSTPRIVATE) + && (array_info = gomp_oacc_needs_data_present (decl))) + { + tree mapping = array_info->mapping; + tree pointer = array_info->pointer; + + 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); + + OMP_CLAUSE_DECL (clause) = unshare_expr (OMP_CLAUSE_DECL (mapping)); + OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_PRESENT); + OMP_CLAUSE_SIZE (clause) = unshare_expr (OMP_CLAUSE_SIZE (mapping)); + + /* Create a new data clause for the firstprivate pointer. */ + tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (clause), + OMP_CLAUSE_MAP); + OMP_CLAUSE_DECL (nc) = unshare_expr (OMP_CLAUSE_DECL (pointer)); + OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER); + + /* For GOMP_MAP_FIRSTPRIVATE_POINTER, this is a bias, not a size. */ + OMP_CLAUSE_SIZE (nc) = unshare_expr (OMP_CLAUSE_SIZE (pointer)); + + /* Create a new data clause for the PSET, if present. */ + tree psetc = NULL; + if (array_info->pset) + { + tree pset = array_info->pset; + psetc = build_omp_clause (OMP_CLAUSE_LOCATION (clause), + OMP_CLAUSE_MAP); + OMP_CLAUSE_DECL (psetc) = unshare_expr (OMP_CLAUSE_DECL (pset)); + OMP_CLAUSE_SIZE (psetc) = unshare_expr (OMP_CLAUSE_SIZE (pset)); + OMP_CLAUSE_SET_MAP_KIND (psetc, GOMP_MAP_TO_PSET); + OMP_CLAUSE_CHAIN (psetc) = nc; + } + + gimplify_omp_ctx *ctx = gimplify_omp_ctxp; + gimplify_omp_ctxp = ctx->outer_context; + gimplify_expr (&OMP_CLAUSE_DECL (clause), pre_p, NULL, + is_gimple_lvalue, fb_lvalue); + gimplify_expr (&OMP_CLAUSE_SIZE (clause), pre_p, NULL, + is_gimple_val, fb_rvalue); + gimplify_expr (&OMP_CLAUSE_SIZE (nc), pre_p, NULL, is_gimple_val, + fb_rvalue); + gimplify_omp_ctxp = ctx; + + OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (clause); + OMP_CLAUSE_CHAIN (clause) = psetc ? psetc : nc; + } 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) diff --git a/gcc/omp-low.c b/gcc/omp-low.c index fdabf67..84c9a88 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -8411,8 +8411,14 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) x = fold_convert_loc (clause_loc, type, x); if (!integer_zerop (OMP_CLAUSE_SIZE (c))) { - tree bias = OMP_CLAUSE_SIZE (c); - if (DECL_P (bias)) + tree bias = OMP_CLAUSE_SIZE (c), remapped_bias; + if (is_gimple_omp_oacc (ctx->stmt)) + { + if (DECL_P (bias) + && (remapped_bias = maybe_lookup_decl (bias, ctx))) + bias = remapped_bias; + } + else if (DECL_P (bias)) bias = lookup_decl (bias, ctx); bias = fold_convert_loc (clause_loc, sizetype, bias); bias = fold_build1_loc (clause_loc, NEGATE_EXPR, sizetype, diff --git a/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c b/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c new file mode 100644 index 0000000..8a039be --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c @@ -0,0 +1,24 @@ +/* Ensure that the gimplifier does not remove any existing clauses as + it inserts new implicit data clauses. */ + +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#define N 100 +static int a[N], b[N]; + +int main(int argc, char *argv[]) +{ + int i; + +#pragma acc data copyin(a[0:N]) copyout (b[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + b[i] = a[i]; + } + + return 0; +} + +// { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:b\\\[0\\\] \\\[len: 400\\\]\\) map\\(to:a\\\[0\\\] \\\[len: 400\\\]\\)" 1 "gimple" } } +// { 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\\\]\\)" 1 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/pr70828.f90 b/gcc/testsuite/gfortran.dg/goacc/pr70828.f90 new file mode 100644 index 0000000..2e58120 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/pr70828.f90 @@ -0,0 +1,22 @@ +! Ensure that pointer mappings are preserved in nested parallel +! constructs. + +! { dg-additional-options "-fdump-tree-gimple" } + +program test + integer, parameter :: n = 100 + integer i, data(n) + + data(:) = 0 + + !$acc data copy(data(5:n-10)) + !$acc parallel loop + do i = 10, n - 10 + data(i) = i + end do + !$acc end parallel loop + !$acc end data +end program test + +! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(tofrom:MEM\\\[\\(c_char \\*\\)\_\[0-9\]+\\\] \\\[len: _\[0-9\]+\\\]\\) map\\(alloc:data \\\[pointer assign, bias: _\[0-9\]+\\\]\\)" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(force_present:MEM\\\[\\(c_char \\*\\)D\\.\[0-9\]+\\\] \\\[len: D\\.\[0-9\]+\\\]\\) map\\(alloc:data \\\[pointer assign, bias: D\\.\[0-9\]+\\\]\\)" 1 "gimple" } } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828-2.c new file mode 100644 index 0000000..357114c --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828-2.c @@ -0,0 +1,34 @@ +/* Subarray declared on data construct, accessed through pointer. */ + +#include <assert.h> + +void +s1 (int *arr, int c) +{ +#pragma acc data copy(arr[5:c-10]) + { +#pragma acc parallel loop + for (int i = 5; i < c - 5; i++) + arr[i] = i; + } +} + +int +main (int argc, char* argv[]) +{ + const int c = 100; + int arr[c]; + + for (int i = 0; i < c; i++) + arr[i] = 0; + + s1 (arr, c); + + for (int i = 0; i < c; i++) + if (i >= 5 && i < c - 5) + assert (arr[i] == i); + else + assert (arr[i] == 0); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c new file mode 100644 index 0000000..4b6dbd7 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c @@ -0,0 +1,27 @@ +/* Subarray declared on enclosing data construct. */ + +#include <assert.h> + +int +main () +{ + int a[100], i; + + for (i = 0; i < 100; i++) + a[i] = 0; + +#pragma acc data copy(a[10:80]) + { + #pragma acc parallel loop + for (i = 10; i < 90; i++) + a[i] = i; + } + + for (i = 0; i < 100; i++) + if (i >= 10 && i < 90) + assert (a[i] == i); + else + assert (a[i] == 0); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/implicit_copy.f90 b/libgomp/testsuite/libgomp.oacc-fortran/implicit_copy.f90 new file mode 100644 index 0000000..7a99f29 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/implicit_copy.f90 @@ -0,0 +1,30 @@ +! { dg-do run } + +integer function test() + implicit none + integer, parameter :: n = 10 + real(8), dimension(n) :: a, b, c + integer i + + do i = 1, n + a(i) = i + b(i) = 1 + end do + + !$acc data copyin(a(1:n), b(1:n)) + !$acc parallel loop + do i = 1, n + c(i) = a(i) * b(i) + end do + !$acc end data + + do i = 1, n + if (c(i) /= a(i) * b(i)) call abort + end do +end function test + +program main + implicit none + integer i, test + i = test() +end program main diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-2.f90 new file mode 100644 index 0000000..22a9566 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-2.f90 @@ -0,0 +1,31 @@ +! Subarrays declared on data construct: assumed-shape array. + +subroutine s1(n, arr) + integer :: n + integer :: arr(n) + + !$acc data copy(arr(5:n-10)) + !$acc parallel loop + do i = 10, n - 10 + arr(i) = i + end do + !$acc end parallel loop + !$acc end data +end subroutine s1 + +program test + integer, parameter :: n = 100 + integer i, data(n) + + data(:) = 0 + + call s1(n, data) + + do i = 1, n + if ((i < 10 .or. i > n-10)) then + if ((data(i) .ne. 0)) call abort + else if (data(i) .ne. i) then + call abort + end if + end do +end program test diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-3.f90 new file mode 100644 index 0000000..ff17d10 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-3.f90 @@ -0,0 +1,34 @@ +! Subarrays declared on data construct: deferred-shape array. + +subroutine s1(n, arr) + integer :: n + integer :: arr(n) + + !$acc data copy(arr(5:n-10)) + !$acc parallel loop + do i = 10, n - 10 + arr(i) = i + end do + !$acc end parallel loop + !$acc end data +end subroutine s1 + +program test + integer, parameter :: n = 100 + integer i + integer, allocatable :: data(:) + + allocate (data(1:n)) + + data(:) = 0 + + call s1(n, data) + + do i = 1, n + if ((i < 10 .or. i > n-10)) then + if ((data(i) .ne. 0)) call abort + else if (data(i) .ne. i) then + call abort + end if + end do +end program test diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90 new file mode 100644 index 0000000..8a16e3d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90 @@ -0,0 +1,29 @@ +! Subarrays on parallel construct (no data construct): assumed-size array. + +subroutine s1(n, arr) + integer :: n + integer :: arr(*) + + !$acc parallel loop copy(arr(5:n-10)) + do i = 10, n - 10 + arr(i) = i + end do + !$acc end parallel loop +end subroutine s1 + +program test + integer, parameter :: n = 100 + integer i, data(n) + + data(:) = 0 + + call s1(n, data) + + do i = 1, n + if ((i < 10 .or. i > n-10)) then + if ((data(i) .ne. 0)) call abort + else if (data(i) .ne. i) then + call abort + end if + end do +end program test diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-6.f90 new file mode 100644 index 0000000..e99c364 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-6.f90 @@ -0,0 +1,28 @@ +! Subarrays declared on data construct: allocatable array (with array +! descriptor). + +program test + integer, parameter :: n = 100 + integer i + integer, allocatable :: data(:) + + allocate (data(1:n)) + + data(:) = 0 + + !$acc data copy(data(5:n-10)) + !$acc parallel loop + do i = 10, n - 10 + data(i) = i + end do + !$acc end parallel loop + !$acc end data + + do i = 1, n + if ((i < 10 .or. i > n-10)) then + if ((data(i) .ne. 0)) call abort + else if (data(i) .ne. i) then + call abort + end if + end do +end program test diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90 new file mode 100644 index 0000000..f87d232 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90 @@ -0,0 +1,24 @@ +! Subarrays on data construct: explicit-shape array. + +program test + integer, parameter :: n = 100 + integer i, data(n) + + data(:) = 0 + + !$acc data copy(data(5:n-10)) + !$acc parallel loop + do i = 10, n - 10 + data(i) = i + end do + !$acc end parallel loop + !$acc end data + + do i = 1, n + if ((i < 10 .or. i > n-10)) then + if ((data(i) .ne. 0)) call abort + else if (data(i) .ne. i) then + call abort + end if + end do +end program test
On Wed, 5 Dec 2018 21:10:45 +0000 Julian Brown <julian@codesourcery.com> wrote: > Thanks for review! How's this version? > > I took the liberty of fixing the patch for Fortran array-descriptor > mappings that use a PSET, also, and adding another test for that > functionality. This is a ping/new version of this patch, incorporating previous review comments and also fixing the inheritance behaviour for references (e.g. for array slices of Fortran function arguments). I've also merged the two patches sent below into one: https://gcc.gnu.org/ml/gcc-patches/2018-08/msg01790.html https://gcc.gnu.org/ml/gcc-patches/2018-08/msg01791.html The second part having been conditionally approved based on approval of the first already, and rebased. Re-tested with offloading to NVPTX. OK? Thanks, Julian 2019-06-09 Julian Brown <julian@codesourcery.com> Cesar Philippidis <cesar@codesourcery.com> gcc/ * gimplify.c (oacc_array_mapping_info): New struct. (gimplify_omp_ctx): Add decl_data_clause hash map. (new_omp_context): Zero-initialise above. (delete_omp_context): Delete above if allocated. (gimplify_scan_omp_clauses): Scan for array mappings on data constructs, and record in above map. (gomp_oacc_needs_data_present): New function. (gimplify_adjust_omp_clauses_1): Handle data mappings (e.g. array slices) declared in lexically-enclosing data constructs. * omp-low.c (lower_omp_target): Allow decl for bias not to be present in OpenACC context. gcc/fortran/ * trans-openmp.c (gfc_omp_finish_clause): Don't raise error for assumed-size array if present in a lexically-enclosing data construct. (gfc_omp_finish_clause): Guard addition of clauses for pointers with DECL_P. gcc/testsuite/ * c-c++-common/goacc/acc-data-chain.c: New test. * gfortran.dg/goacc/pr70828.f90: New test. * gfortran.dg/goacc/pr70828-2.f90: New test. libgomp/ * testsuite/libgomp.oacc-c-c++-common/pr70828.c: New test. * testsuite/libgomp.oacc-fortran/implicit_copy.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828-2.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828-3.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828-4.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828-5.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828-6.f90: New test.
From 9123c4ddd701c40c3e85a0c6cd327066542b9e7a Mon Sep 17 00:00:00 2001 From: Julian Brown <julian@codesourcery.com> Date: Thu, 16 Aug 2018 20:02:10 -0700 Subject: [PATCH 1/2] Inheritance of array sections on data constructs. 2018-08-28 Julian Brown <julian@codesourcery.com> Cesar Philippidis <cesar@codesourcery.com> gcc/ * gimplify.c (gimplify_omp_ctx): Add decl_data_clause hash map. (new_omp_context): Initialise above. (delete_omp_context): Delete above. (gimplify_scan_omp_clauses): Scan for array mappings on data constructs, and record in above map. (gomp_needs_data_present): New function. (gimplify_adjust_omp_clauses_1): Handle data mappings (e.g. array slices) declared in lexically-enclosing data constructs. * omp-low.c (lower_omp_target): Allow decl for bias not to be present in omp context. gcc/testsuite/ * c-c++-common/goacc/acc-data-chain.c: New test. * gfortran.dg/goacc/pr70828.f90: New test. * gfortran.dg/goacc/pr70828-2.f90: New test. libgomp/ * testsuite/libgomp.oacc-c-c++-common/pr70828.c: New test. * testsuite/libgomp.oacc-fortran/implicit_copy.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828-2.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828-3.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828-5.f90: New test. --- gcc/gimplify.c | 97 +++++++++++++++++++++- gcc/omp-low.c | 7 +- gcc/testsuite/c-c++-common/goacc/acc-data-chain.c | 24 ++++++ gcc/testsuite/gfortran.dg/goacc/pr70828.f90 | 22 +++++ .../libgomp.oacc-c-c++-common/pr70828-2.c | 34 ++++++++ .../testsuite/libgomp.oacc-c-c++-common/pr70828.c | 27 ++++++ .../libgomp.oacc-fortran/implicit_copy.f90 | 30 +++++++ .../testsuite/libgomp.oacc-fortran/pr70828-2.f90 | 31 +++++++ .../testsuite/libgomp.oacc-fortran/pr70828-3.f90 | 34 ++++++++ .../testsuite/libgomp.oacc-fortran/pr70828-5.f90 | 29 +++++++ libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90 | 24 ++++++ 11 files changed, 354 insertions(+), 5 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/acc-data-chain.c create mode 100644 gcc/testsuite/gfortran.dg/goacc/pr70828.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/implicit_copy.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/pr70828-2.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/pr70828-3.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90 diff --git a/gcc/gimplify.c b/gcc/gimplify.c index dbd0f0e..d704aef 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -191,6 +191,7 @@ struct gimplify_omp_ctx bool target_map_scalars_firstprivate; bool target_map_pointers_as_0len_arrays; bool target_firstprivatize_array_bases; + hash_map<tree, std::pair<tree, tree> > *decl_data_clause; }; static struct gimplify_ctx *gimplify_ctxp; @@ -413,6 +414,7 @@ new_omp_context (enum omp_region_type region_type) c->default_kind = OMP_CLAUSE_DEFAULT_SHARED; else c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED; + c->decl_data_clause = new hash_map<tree, std::pair<tree, tree> >; return c; } @@ -425,6 +427,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); } @@ -7793,8 +7796,21 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_TARGET: break; case OACC_DATA: - if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE) - break; + { + tree nextc = OMP_CLAUSE_CHAIN (c); + if (nextc + && OMP_CLAUSE_CODE (nextc) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (nextc) + == GOMP_MAP_FIRSTPRIVATE_POINTER + || OMP_CLAUSE_MAP_KIND (nextc) == GOMP_MAP_POINTER)) + { + tree base_addr = OMP_CLAUSE_DECL (nextc); + ctx->decl_data_clause->put (base_addr, + std::make_pair (unshare_expr (c), unshare_expr (nextc))); + } + if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE) + break; + } /* FALLTHRU */ case OMP_TARGET_DATA: case OMP_TARGET_ENTER_DATA: @@ -8692,6 +8708,45 @@ struct gimplify_adjust_omp_clauses_data gimple_seq *pre_p; }; +/* For OpenACC parallel and kernels 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 parallel or + kernels region. This function returns a pair consisting of the mapping for + the data itself and for the pointer to the beginning of the data if present + in an outer data construct, or returns NULL otherwise. */ + +static std::pair<tree, tree> * +gomp_needs_data_present (tree decl) +{ + gimplify_omp_ctx *ctx = NULL; + + if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE + && TREE_CODE (TREE_TYPE (decl)) != POINTER_TYPE + && (TREE_CODE (TREE_TYPE (decl)) != POINTER_TYPE + || TREE_CODE (TREE_TYPE (TREE_TYPE (decl))) != ARRAY_TYPE)) + return NULL; + + if (gimplify_omp_ctxp->region_type != ORT_ACC_PARALLEL + && gimplify_omp_ctxp->region_type != ORT_ACC_KERNELS) + return NULL; + + decl = get_base_address (decl); + + for (ctx = gimplify_omp_ctxp->outer_context; ctx; ctx = ctx->outer_context) + { + std::pair<tree, tree> *ret; + + if (ctx->region_type != ORT_ACC_DATA) + break; + + if ((ret = ctx->decl_data_clause->get (decl))) + return ret; + } + + return NULL; +} + /* For all variables that were not actually used within the context, remove PRIVATE, SHARED, and FIRSTPRIVATE clauses. */ @@ -8784,6 +8839,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; + std::pair<tree, tree> *mapping_ptr; if (private_debug) OMP_CLAUSE_PRIVATE_DEBUG (clause) = 1; else if (code == OMP_CLAUSE_PRIVATE && (flags & GOVD_PRIVATE_OUTER_REF)) @@ -8792,6 +8848,43 @@ 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 ((code == OMP_CLAUSE_MAP || code == OMP_CLAUSE_FIRSTPRIVATE) + && (mapping_ptr = gomp_needs_data_present (decl))) + { + tree mapping = mapping_ptr->first; + tree pointer = mapping_ptr->second; + + 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); + + OMP_CLAUSE_DECL (clause) = unshare_expr (OMP_CLAUSE_DECL (mapping)); + OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_PRESENT); + OMP_CLAUSE_SIZE (clause) = unshare_expr (OMP_CLAUSE_SIZE (mapping)); + + /* Create a new data clause for the firstprivate pointer. */ + tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (clause), + OMP_CLAUSE_MAP); + OMP_CLAUSE_DECL (nc) = unshare_expr (OMP_CLAUSE_DECL (pointer)); + OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER); + + /* For GOMP_MAP_FIRSTPRIVATE_POINTER, this is a bias, not a size. */ + OMP_CLAUSE_SIZE (nc) = unshare_expr (OMP_CLAUSE_SIZE (pointer)); + + gimplify_omp_ctx *ctx = gimplify_omp_ctxp; + gimplify_omp_ctxp = ctx->outer_context; + gimplify_expr (&OMP_CLAUSE_DECL (clause), pre_p, NULL, + is_gimple_lvalue, fb_lvalue); + gimplify_expr (&OMP_CLAUSE_SIZE (clause), pre_p, NULL, + is_gimple_val, fb_rvalue); + gimplify_expr (&OMP_CLAUSE_SIZE (nc), pre_p, NULL, is_gimple_val, + fb_rvalue); + gimplify_omp_ctxp = ctx; + + OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (clause); + OMP_CLAUSE_CHAIN (clause) = nc; + } 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) diff --git a/gcc/omp-low.c b/gcc/omp-low.c index fdabf67..be2bb73 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -8411,9 +8411,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) x = fold_convert_loc (clause_loc, type, x); if (!integer_zerop (OMP_CLAUSE_SIZE (c))) { - tree bias = OMP_CLAUSE_SIZE (c); - if (DECL_P (bias)) - bias = lookup_decl (bias, ctx); + tree bias = OMP_CLAUSE_SIZE (c), remapped_bias; + if (DECL_P (bias) + && (remapped_bias = maybe_lookup_decl (bias, ctx))) + bias = remapped_bias; bias = fold_convert_loc (clause_loc, sizetype, bias); bias = fold_build1_loc (clause_loc, NEGATE_EXPR, sizetype, bias); diff --git a/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c b/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c new file mode 100644 index 0000000..8a039be --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c @@ -0,0 +1,24 @@ +/* Ensure that the gimplifier does not remove any existing clauses as + it inserts new implicit data clauses. */ + +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#define N 100 +static int a[N], b[N]; + +int main(int argc, char *argv[]) +{ + int i; + +#pragma acc data copyin(a[0:N]) copyout (b[0:N]) + { +#pragma acc parallel loop + for (i = 0; i < N; i++) + b[i] = a[i]; + } + + return 0; +} + +// { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:b\\\[0\\\] \\\[len: 400\\\]\\) map\\(to:a\\\[0\\\] \\\[len: 400\\\]\\)" 1 "gimple" } } +// { 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\\\]\\)" 1 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/pr70828.f90 b/gcc/testsuite/gfortran.dg/goacc/pr70828.f90 new file mode 100644 index 0000000..2e58120 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/pr70828.f90 @@ -0,0 +1,22 @@ +! Ensure that pointer mappings are preserved in nested parallel +! constructs. + +! { dg-additional-options "-fdump-tree-gimple" } + +program test + integer, parameter :: n = 100 + integer i, data(n) + + data(:) = 0 + + !$acc data copy(data(5:n-10)) + !$acc parallel loop + do i = 10, n - 10 + data(i) = i + end do + !$acc end parallel loop + !$acc end data +end program test + +! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(tofrom:MEM\\\[\\(c_char \\*\\)\_\[0-9\]+\\\] \\\[len: _\[0-9\]+\\\]\\) map\\(alloc:data \\\[pointer assign, bias: _\[0-9\]+\\\]\\)" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(force_present:MEM\\\[\\(c_char \\*\\)D\\.\[0-9\]+\\\] \\\[len: D\\.\[0-9\]+\\\]\\) map\\(alloc:data \\\[pointer assign, bias: D\\.\[0-9\]+\\\]\\)" 1 "gimple" } } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828-2.c new file mode 100644 index 0000000..357114c --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828-2.c @@ -0,0 +1,34 @@ +/* Subarray declared on data construct, accessed through pointer. */ + +#include <assert.h> + +void +s1 (int *arr, int c) +{ +#pragma acc data copy(arr[5:c-10]) + { +#pragma acc parallel loop + for (int i = 5; i < c - 5; i++) + arr[i] = i; + } +} + +int +main (int argc, char* argv[]) +{ + const int c = 100; + int arr[c]; + + for (int i = 0; i < c; i++) + arr[i] = 0; + + s1 (arr, c); + + for (int i = 0; i < c; i++) + if (i >= 5 && i < c - 5) + assert (arr[i] == i); + else + assert (arr[i] == 0); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c new file mode 100644 index 0000000..4b6dbd7 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c @@ -0,0 +1,27 @@ +/* Subarray declared on enclosing data construct. */ + +#include <assert.h> + +int +main () +{ + int a[100], i; + + for (i = 0; i < 100; i++) + a[i] = 0; + +#pragma acc data copy(a[10:80]) + { + #pragma acc parallel loop + for (i = 10; i < 90; i++) + a[i] = i; + } + + for (i = 0; i < 100; i++) + if (i >= 10 && i < 90) + assert (a[i] == i); + else + assert (a[i] == 0); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/implicit_copy.f90 b/libgomp/testsuite/libgomp.oacc-fortran/implicit_copy.f90 new file mode 100644 index 0000000..7a99f29 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/implicit_copy.f90 @@ -0,0 +1,30 @@ +! { dg-do run } + +integer function test() + implicit none + integer, parameter :: n = 10 + real(8), dimension(n) :: a, b, c + integer i + + do i = 1, n + a(i) = i + b(i) = 1 + end do + + !$acc data copyin(a(1:n), b(1:n)) + !$acc parallel loop + do i = 1, n + c(i) = a(i) * b(i) + end do + !$acc end data + + do i = 1, n + if (c(i) /= a(i) * b(i)) call abort + end do +end function test + +program main + implicit none + integer i, test + i = test() +end program main diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-2.f90 new file mode 100644 index 0000000..22a9566 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-2.f90 @@ -0,0 +1,31 @@ +! Subarrays declared on data construct: assumed-shape array. + +subroutine s1(n, arr) + integer :: n + integer :: arr(n) + + !$acc data copy(arr(5:n-10)) + !$acc parallel loop + do i = 10, n - 10 + arr(i) = i + end do + !$acc end parallel loop + !$acc end data +end subroutine s1 + +program test + integer, parameter :: n = 100 + integer i, data(n) + + data(:) = 0 + + call s1(n, data) + + do i = 1, n + if ((i < 10 .or. i > n-10)) then + if ((data(i) .ne. 0)) call abort + else if (data(i) .ne. i) then + call abort + end if + end do +end program test diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-3.f90 new file mode 100644 index 0000000..ff17d10 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-3.f90 @@ -0,0 +1,34 @@ +! Subarrays declared on data construct: deferred-shape array. + +subroutine s1(n, arr) + integer :: n + integer :: arr(n) + + !$acc data copy(arr(5:n-10)) + !$acc parallel loop + do i = 10, n - 10 + arr(i) = i + end do + !$acc end parallel loop + !$acc end data +end subroutine s1 + +program test + integer, parameter :: n = 100 + integer i + integer, allocatable :: data(:) + + allocate (data(1:n)) + + data(:) = 0 + + call s1(n, data) + + do i = 1, n + if ((i < 10 .or. i > n-10)) then + if ((data(i) .ne. 0)) call abort + else if (data(i) .ne. i) then + call abort + end if + end do +end program test diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90 new file mode 100644 index 0000000..8a16e3d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90 @@ -0,0 +1,29 @@ +! Subarrays on parallel construct (no data construct): assumed-size array. + +subroutine s1(n, arr) + integer :: n + integer :: arr(*) + + !$acc parallel loop copy(arr(5:n-10)) + do i = 10, n - 10 + arr(i) = i + end do + !$acc end parallel loop +end subroutine s1 + +program test + integer, parameter :: n = 100 + integer i, data(n) + + data(:) = 0 + + call s1(n, data) + + do i = 1, n + if ((i < 10 .or. i > n-10)) then + if ((data(i) .ne. 0)) call abort + else if (data(i) .ne. i) then + call abort + end if + end do +end program test diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90 new file mode 100644 index 0000000..f87d232 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90 @@ -0,0 +1,24 @@ +! Subarrays on data construct: explicit-shape array. + +program test + integer, parameter :: n = 100 + integer i, data(n) + + data(:) = 0 + + !$acc data copy(data(5:n-10)) + !$acc parallel loop + do i = 10, n - 10 + data(i) = i + end do + !$acc end parallel loop + !$acc end data + + do i = 1, n + if ((i < 10 .or. i > n-10)) then + if ((data(i) .ne. 0)) call abort + else if (data(i) .ne. i) then + call abort + end if + end do +end program test -- 1.8.1.1
This patch implements support for array slices (with a non-zero base element) declared on OpenACC data constructs. Any lexically-enclosed parallel or kernels regions should "inherit" such mappings, e.g. if we have: #pragma acc data copy(arr[10:20]) { #pragma acc parallel loop for (...) { ...arr[X]... } } the mapping for "arr" on the data construct takes precedence over the default mapping behaviour for the parallel construct, which is to map the whole array. (OpenACC 2.5, "2.5.1. Parallel Construct" and elsewhere). Tested with offloading to nvptx. (This patch differs in implementation somewhat from the version on the gomp4, etc. branches.) OK to apply? Thanks, Julian 2018-08-28 Julian Brown <julian@codesourcery.com> Cesar Philippidis <cesar@codesourcery.com> PR middle-end/70828 gcc/ * gimplify.c (gimplify_omp_ctx): Add decl_data_clause hash map. (new_omp_context): Initialise above. (delete_omp_context): Delete above. (gimplify_scan_omp_clauses): Scan for array mappings on data constructs, and record in above map. (gomp_needs_data_present): New function. (gimplify_adjust_omp_clauses_1): Handle data mappings (e.g. array slices) declared in lexically-enclosing data constructs. * omp-low.c (lower_omp_target): Allow decl for bias not to be present in omp context. gcc/testsuite/ * c-c++-common/goacc/acc-data-chain.c: New test. * gfortran.dg/goacc/pr70828.f90: New test. * gfortran.dg/goacc/pr70828-2.f90: New test. libgomp/ * testsuite/libgomp.oacc-c-c++-common/pr70828.c: New test. * testsuite/libgomp.oacc-fortran/implicit_copy.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828-2.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828-3.f90: New test. * testsuite/libgomp.oacc-fortran/pr70828-5.f90: New test.