commit 390adf97cfdde951ed1e82fc54d77e34130c70b8 Author: Julian Brown Date: Thu Aug 16 20:02:10 2018 -0700 Inheritance of array sections on data constructs. 2018-08-28 Julian Brown Cesar Philippidis 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 *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; + 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 + +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 + +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