From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 90276 invoked by alias); 28 Aug 2018 19:19:38 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 90264 invoked by uid 89); 28 Aug 2018 19:19:37 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-24.4 required=5.0 tests=AWL,BAYES_00,GIT_PATCH_0,GIT_PATCH_1,GIT_PATCH_2,GIT_PATCH_3,RCVD_IN_DNSWL_NONE,SPF_PASS,URIBL_RED autolearn=ham version=3.3.2 spammy=PRIVATE, brown, Brown X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Tue, 28 Aug 2018 19:19:34 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-MBX-04.mgc.mentorg.com) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1fujWV-0000jd-FB from Julian_Brown@mentor.com ; Tue, 28 Aug 2018 12:19:31 -0700 Received: from squid.athome (137.202.0.90) by SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Tue, 28 Aug 2018 20:19:26 +0100 Date: Tue, 28 Aug 2018 19:19:00 -0000 From: Julian Brown To: , , Jakub Jelinek Subject: [PATCH, OpenACC] (1/2) Fix implicit mapping for array slices on lexically-enclosing data constructs (PR70828) Message-ID: <20180828151919.576c636c@squid.athome> MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="MP_/QFroXZjOKvMcvuAsiB8y+iG" X-IsSubscribed: yes X-SW-Source: 2018-08/txt/msg01790.txt.bz2 --MP_/QFroXZjOKvMcvuAsiB8y+iG Content-Type: text/plain; charset="US-ASCII" Content-Transfer-Encoding: 7bit Content-Disposition: inline Content-length: 1836 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 Cesar Philippidis 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. --MP_/QFroXZjOKvMcvuAsiB8y+iG Content-Type: text/x-patch Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="0001-Inheritance-of-array-sections-on-data-constructs.patch" Content-length: 16944 >From 9123c4ddd701c40c3e85a0c6cd327066542b9e7a Mon Sep 17 00:00:00 2001 From: Julian Brown 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 Cesar Philippidis 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 > *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 >; 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 * +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 *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 *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 + +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.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 --MP_/QFroXZjOKvMcvuAsiB8y+iG--