From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1729) id 4B06F3896C17; Thu, 13 May 2021 16:09:01 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 4B06F3896C17 Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Kwok Yeung To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-11] Add OpenACC Fortran support for deviceptr and variable in common blocks X-Act-Checkin: gcc X-Git-Author: Julian Brown X-Git-Refname: refs/heads/devel/omp/gcc-11 X-Git-Oldrev: 5c357d4f2c835f9b4f432ab03847beb11aec8888 X-Git-Newrev: a14b3f29681da1d2465e15f98b8cf8d5c64a2c3c Message-Id: <20210513160901.4B06F3896C17@sourceware.org> Date: Thu, 13 May 2021 16:09:01 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Thu, 13 May 2021 16:09:01 -0000 https://gcc.gnu.org/g:a14b3f29681da1d2465e15f98b8cf8d5c64a2c3c commit a14b3f29681da1d2465e15f98b8cf8d5c64a2c3c Author: Julian Brown Date: Tue Feb 12 14:32:34 2019 -0800 Add OpenACC Fortran support for deviceptr and variable in common blocks 2018-06-29 Cesar Philippidis James Norris gcc/fortran/ * openmp.c (gfc_match_omp_map_clause): Re-write handling of the deviceptr clause. Add new common_blocks argument. Propagate it to gfc_match_omp_variable_list. (gfc_match_omp_clauses): Update calls to gfc_match_omp_map_clauses. (resolve_positive_int_expr): Promote the warning to an error. (check_array_not_assumed): Remove pointer check. (resolve_oacc_nested_loops): Error on do concurrent loops. * trans-openmp.c (gfc_omp_finish_clause): Don't create pointer data mappings for deviceptr clauses. (gfc_trans_omp_clauses): Likewise. gcc/ * gimplify.c (enum gimplify_omp_var_data): Add GOVD_DEVICETPR. (oacc_default_clause): Privatize fortran common blocks. (omp_notice_variable): Add GOVD_DEVICEPTR attribute when appropriate. Defer the expansion of DECL_VALUE_EXPR for common block decls. (gimplify_scan_omp_clauses): Add GOVD_DEVICEPTR attribute when appropriate. (gimplify_adjust_omp_clauses_1): Set GOMP_MAP_FORCE_DEVICEPTR for implicit deviceptr mappings. gcc/testsuite/ * c-c++-common/goacc/deviceptr-4.c: Update. * gfortran.dg/goacc/common-block-1.f90: New test. * gfortran.dg/goacc/common-block-2.f90: New test. * gfortran.dg/goacc/loop-2.f95: Update. * gfortran.dg/goacc/loop-3-2.f95: Update. * gfortran.dg/goacc/loop-3.f95: Update. * gfortran.dg/goacc/pr72715.f90: New test. * gfortran.dg/goacc/sie.f95: Update. * gfortran.dg/goacc/tile-1.f90: Update. * gfortran.dg/gomp/pr77516.f90: Update. libgomp/ * oacc-parallel.c (GOACC_parallel_keyed): Handle Fortran deviceptr clause. (GOACC_data_start): Likewise. * testsuite/libgomp.oacc-fortran/common-block-1.f90: New test. * testsuite/libgomp.oacc-fortran/common-block-2.f90: New test. * testsuite/libgomp.oacc-fortran/common-block-3.f90: New test. * testsuite/libgomp.oacc-fortran/deviceptr-1.f90: New test. Diff: --- gcc/ChangeLog.omp | 10 ++ gcc/fortran/ChangeLog.omp | 6 + gcc/fortran/openmp.c | 4 +- gcc/fortran/trans-openmp.c | 9 + gcc/gimplify.c | 12 +- gcc/testsuite/ChangeLog.omp | 12 ++ gcc/testsuite/c-c++-common/goacc/deviceptr-4.c | 2 +- .../gfortran.dg/goacc/loop-2-kernels-tile.f95 | 4 +- .../gfortran.dg/goacc/loop-2-parallel-tile.f95 | 4 +- gcc/testsuite/gfortran.dg/goacc/sie.f95 | 36 ++-- gcc/testsuite/gfortran.dg/goacc/tile-1.f90 | 16 +- gcc/testsuite/gfortran.dg/gomp/pr77516.f90 | 2 +- libgomp/ChangeLog.omp | 8 + libgomp/oacc-parallel.c | 2 + .../testsuite/libgomp.oacc-fortran/deviceptr-1.f90 | 197 +++++++++++++++++++++ 15 files changed, 289 insertions(+), 35 deletions(-) diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index b3886838368..a953a59af9a 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,13 @@ +2018-06-29 Cesar Philippidis + James Norris + + * gimplify.c (enum gimplify_omp_var_data): Add GOVD_DEVICETPR. + (omp_notice_variable): Add GOVD_DEVICEPTR attribute when appropriate. + (gimplify_scan_omp_clauses): Add GOVD_DEVICEPTR attribute when + appropriate. + (gimplify_adjust_omp_clauses_1): Set GOMP_MAP_FORCE_DEVICEPTR for + implicit deviceptr mappings. + 2018-10-02 Thomas Schwinge Cesar Philippidis diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp index 24f67c1568d..5020b76e682 100644 --- a/gcc/fortran/ChangeLog.omp +++ b/gcc/fortran/ChangeLog.omp @@ -1,3 +1,9 @@ +2018-06-29 Cesar Philippidis + James Norris + + * openmp.c (resolve_positive_int_expr): Promote the warning to an + error. + 2018-10-02 Thomas Schwinge Cesar Philippidis diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 068cf79307e..d8186117f1a 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -4504,8 +4504,8 @@ resolve_positive_int_expr (gfc_expr *expr, const char *clause) if (expr->expr_type == EXPR_CONSTANT && expr->ts.type == BT_INTEGER && mpz_sgn (expr->value.integer) <= 0) - gfc_warning (0, "INTEGER expression of %s clause at %L must be positive", - clause, &expr->where); + gfc_error ("INTEGER expression of %s clause at %L must be positive", + clause, &expr->where); } static void diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index c5c8e3f0cee..689649e50e2 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -1306,6 +1306,9 @@ gfc_omp_finish_clause (tree c, gimple_seq *pre_p, bool openacc) return; } + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR) + return; + tree c2 = NULL_TREE, c3 = NULL_TREE, c4 = NULL_TREE; tree present = gfc_omp_check_optional_argument (decl, true); if (POINTER_TYPE_P (TREE_TYPE (decl))) @@ -2736,6 +2739,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, OMP_CLAUSE_SIZE (node3) = size_int (0); goto finalize_map_clause; } + else if (POINTER_TYPE_P (TREE_TYPE (decl)) + && n->u.map_op == OMP_MAP_FORCE_DEVICEPTR) + { + OMP_CLAUSE_DECL (node) = decl; + goto finalize_map_clause; + } else if (POINTER_TYPE_P (TREE_TYPE (decl)) && (gfc_omp_privatize_by_reference (decl) || GFC_DECL_GET_SCALAR_POINTER (decl) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index fa7a002e584..d5c7d8e1c0e 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -128,6 +128,9 @@ enum gimplify_omp_var_data fields. */ GOVD_MAP_HAS_ATTACHMENTS = 8388608, + /* Flag for OpenACC deviceptrs. */ + GOVD_DEVICEPTR = (1<<24), + GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR | GOVD_LOCAL) @@ -7574,6 +7577,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) error ("variable %qE declared in enclosing " "% region", DECL_NAME (decl)); nflags |= GOVD_MAP; + nflags |= (n2->value & GOVD_DEVICEPTR); if (octx->region_type == ORT_ACC_DATA && (n2->value & GOVD_MAP_0LEN_ARRAY)) nflags |= GOVD_MAP_0LEN_ARRAY; @@ -9484,6 +9488,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM) flags |= GOVD_MAP_ALWAYS_TO; + else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR) + flags |= GOVD_DEVICEPTR; if ((code == OMP_TARGET || code == OMP_TARGET_DATA @@ -10356,7 +10362,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) | GOVD_MAP_FORCE | GOVD_MAP_FORCE_PRESENT | GOVD_MAP_ALLOC_ONLY - | GOVD_MAP_FROM_ONLY)) + | GOVD_MAP_FROM_ONLY + | GOVD_DEVICEPTR)) { case 0: kind = GOMP_MAP_TOFROM; @@ -10379,6 +10386,9 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) case GOVD_MAP_FORCE_PRESENT: kind = GOMP_MAP_FORCE_PRESENT; break; + case GOVD_DEVICEPTR: + kind = GOMP_MAP_FORCE_DEVICEPTR; + break; default: gcc_unreachable (); } diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index 0854ba050ba..5d1317d7941 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,15 @@ +2018-06-29 Cesar Philippidis + James Norris + + * c-c++-common/goacc/deviceptr-4.c: Update. + * gfortran.dg/goacc/loop-2.f95: Update. + * gfortran.dg/goacc/loop-3-2.f95: Update. + * gfortran.dg/goacc/loop-3.f95: Update. + * gfortran.dg/goacc/pr72715.f90: New test. + * gfortran.dg/goacc/sie.f95: Update. + * gfortran.dg/goacc/tile-1.f90: Update. + * gfortran.dg/gomp/pr77516.f90: Update. + 2018-10-02 Thomas Schwinge Cesar Philippidis diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c index db1b91633a6..79a51620db9 100644 --- a/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c +++ b/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c @@ -8,4 +8,4 @@ subr (int *a) a[0] += 1.0; } -/* { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel.*map\\(tofrom:a" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel.*map\\(force_deviceptr:a" 1 "gimple" } } */ diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels-tile.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels-tile.f95 index afc8a278cac..65425159a2c 100644 --- a/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels-tile.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels-tile.f95 @@ -29,7 +29,7 @@ program test DO j = 1,10 ENDDO ENDDO - !$acc loop tile(-1) ! { dg-warning "must be positive" } + !$acc loop tile(-1) ! { dg-error "must be positive" } do i = 1,10 enddo !$acc loop tile(i) ! { dg-error "constant expression" } @@ -82,7 +82,7 @@ program test DO j = 1,10 ENDDO ENDDO - !$acc kernels loop tile(-1) ! { dg-warning "must be positive" } + !$acc kernels loop tile(-1) ! { dg-error "must be positive" } do i = 1,10 enddo !$acc kernels loop tile(i) ! { dg-error "constant expression" } diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-2-parallel-tile.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-2-parallel-tile.f95 index 4bfca748f75..dae8f667486 100644 --- a/gcc/testsuite/gfortran.dg/goacc/loop-2-parallel-tile.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/loop-2-parallel-tile.f95 @@ -20,7 +20,7 @@ program test DO j = 1,10 ENDDO ENDDO - !$acc loop tile(-1) ! { dg-warning "must be positive" } + !$acc loop tile(-1) ! { dg-error "must be positive" } do i = 1,10 enddo !$acc loop tile(i) ! { dg-error "constant expression" } @@ -73,7 +73,7 @@ program test DO j = 1,10 ENDDO ENDDO - !$acc parallel loop tile(-1) ! { dg-warning "must be positive" } + !$acc parallel loop tile(-1) ! { dg-error "must be positive" } do i = 1,10 enddo !$acc parallel loop tile(i) ! { dg-error "constant expression" } diff --git a/gcc/testsuite/gfortran.dg/goacc/sie.f95 b/gcc/testsuite/gfortran.dg/goacc/sie.f95 index 194a1daae5f..bb942bdb533 100644 --- a/gcc/testsuite/gfortran.dg/goacc/sie.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/sie.f95 @@ -78,10 +78,10 @@ program test !$acc parallel num_gangs(i+1) !$acc end parallel - !$acc parallel num_gangs(-1) ! { dg-warning "must be positive" } + !$acc parallel num_gangs(-1) ! { dg-error "must be positive" } !$acc end parallel - !$acc parallel num_gangs(0) ! { dg-warning "must be positive" } + !$acc parallel num_gangs(0) ! { dg-error "must be positive" } !$acc end parallel !$acc parallel num_gangs() ! { dg-error "Invalid character in name" } @@ -106,10 +106,10 @@ program test !$acc kernels num_gangs(i+1) !$acc end kernels - !$acc kernels num_gangs(-1) ! { dg-warning "must be positive" } + !$acc kernels num_gangs(-1) ! { dg-error "must be positive" } !$acc end kernels - !$acc kernels num_gangs(0) ! { dg-warning "must be positive" } + !$acc kernels num_gangs(0) ! { dg-error "must be positive" } !$acc end kernels !$acc kernels num_gangs() ! { dg-error "Invalid character in name" } @@ -135,10 +135,10 @@ program test !$acc parallel num_workers(i+1) !$acc end parallel - !$acc parallel num_workers(-1) ! { dg-warning "must be positive" } + !$acc parallel num_workers(-1) ! { dg-error "must be positive" } !$acc end parallel - !$acc parallel num_workers(0) ! { dg-warning "must be positive" } + !$acc parallel num_workers(0) ! { dg-error "must be positive" } !$acc end parallel !$acc parallel num_workers() ! { dg-error "Invalid character in name" } @@ -163,10 +163,10 @@ program test !$acc kernels num_workers(i+1) !$acc end kernels - !$acc kernels num_workers(-1) ! { dg-warning "must be positive" } + !$acc kernels num_workers(-1) ! { dg-error "must be positive" } !$acc end kernels - !$acc kernels num_workers(0) ! { dg-warning "must be positive" } + !$acc kernels num_workers(0) ! { dg-error "must be positive" } !$acc end kernels !$acc kernels num_workers() ! { dg-error "Invalid character in name" } @@ -192,10 +192,10 @@ program test !$acc parallel vector_length(i+1) !$acc end parallel - !$acc parallel vector_length(-1) ! { dg-warning "must be positive" } + !$acc parallel vector_length(-1) ! { dg-error "must be positive" } !$acc end parallel - !$acc parallel vector_length(0) ! { dg-warning "must be positive" } + !$acc parallel vector_length(0) ! { dg-error "must be positive" } !$acc end parallel !$acc parallel vector_length() ! { dg-error "Invalid character in name" } @@ -220,10 +220,10 @@ program test !$acc kernels vector_length(i+1) !$acc end kernels - !$acc kernels vector_length(-1) ! { dg-warning "must be positive" } + !$acc kernels vector_length(-1) ! { dg-error "must be positive" } !$acc end kernels - !$acc kernels vector_length(0) ! { dg-warning "must be positive" } + !$acc kernels vector_length(0) ! { dg-error "must be positive" } !$acc end kernels !$acc kernels vector_length() ! { dg-error "Invalid character in name" } @@ -250,10 +250,10 @@ program test !$acc loop gang(i+1) do i = 1,10 enddo - !$acc loop gang(-1) ! { dg-warning "must be positive" } + !$acc loop gang(-1) ! { dg-error "must be positive" } do i = 1,10 enddo - !$acc loop gang(0) ! { dg-warning "must be positive" } + !$acc loop gang(0) ! { dg-error "must be positive" } do i = 1,10 enddo !$acc loop gang() ! { dg-error "Invalid character in name" } @@ -282,10 +282,10 @@ program test !$acc loop worker(i+1) do i = 1,10 enddo - !$acc loop worker(-1) ! { dg-warning "must be positive" } + !$acc loop worker(-1) ! { dg-error "must be positive" } do i = 1,10 enddo - !$acc loop worker(0) ! { dg-warning "must be positive" } + !$acc loop worker(0) ! { dg-error "must be positive" } do i = 1,10 enddo !$acc loop worker() ! { dg-error "Invalid character in name" } @@ -314,10 +314,10 @@ program test !$acc loop vector(i+1) do i = 1,10 enddo - !$acc loop vector(-1) ! { dg-warning "must be positive" } + !$acc loop vector(-1) ! { dg-error "must be positive" } do i = 1,10 enddo - !$acc loop vector(0) ! { dg-warning "must be positive" } + !$acc loop vector(0) ! { dg-error "must be positive" } do i = 1,10 enddo !$acc loop vector() ! { dg-error "Invalid character in name" } diff --git a/gcc/testsuite/gfortran.dg/goacc/tile-1.f90 b/gcc/testsuite/gfortran.dg/goacc/tile-1.f90 index f609b127df9..9ef75211087 100644 --- a/gcc/testsuite/gfortran.dg/goacc/tile-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/tile-1.f90 @@ -44,17 +44,17 @@ subroutine parloop do i = 1, n end do - !$acc parallel loop tile(-3) ! { dg-warning "must be positive" } + !$acc parallel loop tile(-3) ! { dg-error "must be positive" } do i = 1, n end do - !$acc parallel loop tile(10, -3) ! { dg-warning "must be positive" } + !$acc parallel loop tile(10, -3) ! { dg-error "must be positive" } do i = 1, n do j = 1, n end do end do - !$acc parallel loop tile(-100, 10, 5) ! { dg-warning "must be positive" } + !$acc parallel loop tile(-100, 10, 5) ! { dg-error "must be positive" } do i = 1, n do j = 1, n do k = 1, n @@ -114,7 +114,7 @@ subroutine par end do end do - !$acc loop tile(-2) ! { dg-warning "must be positive" } + !$acc loop tile(-2) ! { dg-error "must be positive" } do i = 1, n end do @@ -195,7 +195,7 @@ subroutine kern end do end do - !$acc loop tile(-2) ! { dg-warning "must be positive" } + !$acc loop tile(-2) ! { dg-error "must be positive" } do i = 1, n end do @@ -295,17 +295,17 @@ subroutine kernsloop do i = 1, n end do - !$acc kernels loop tile(-3) ! { dg-warning "must be positive" } + !$acc kernels loop tile(-3) ! { dg-error "must be positive" } do i = 1, n end do - !$acc kernels loop tile(10, -3) ! { dg-warning "must be positive" } + !$acc kernels loop tile(10, -3) ! { dg-error "must be positive" } do i = 1, n do j = 1, n end do end do - !$acc kernels loop tile(-100, 10, 5) ! { dg-warning "must be positive" } + !$acc kernels loop tile(-100, 10, 5) ! { dg-error "must be positive" } do i = 1, n do j = 1, n do k = 1, n diff --git a/gcc/testsuite/gfortran.dg/gomp/pr77516.f90 b/gcc/testsuite/gfortran.dg/gomp/pr77516.f90 index 9c0a95b9f79..3ac3f5562d0 100644 --- a/gcc/testsuite/gfortran.dg/gomp/pr77516.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/pr77516.f90 @@ -4,7 +4,7 @@ program pr77516 integer :: i, x x = 0 -!$omp simd safelen(0) reduction(+:x) ! { dg-warning "must be positive" } +!$omp simd safelen(0) reduction(+:x) ! { dg-error "must be positive" } do i = 1, 8 x = x + 1 end do diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index a18e49aa572..df624ebfd41 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,11 @@ +2018-06-29 Cesar Philippidis + James Norris + + * oacc-parallel.c (GOACC_parallel_keyed): Handle Fortran deviceptr + clause. + (GOACC_data_start): Likewise. + * testsuite/libgomp.oacc-fortran/deviceptr-1.f90: New test. + 2019-02-12 Julian Brown * oacc-cuda.c (acc_set_cuda_stream): Return 0 on error/invalid diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index 55e16427e73..0afab0b29ae 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -643,6 +643,8 @@ GOACC_data_start (int flags_m, size_t mapnum, if (profiling_p) goacc_profiling_dispatch (&prof_info, &enter_data_event_info, &api_info); + handle_ftn_pointers (mapnum, hostaddrs, sizes, kinds); + /* Host fallback or 'do nothing'. */ if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) || (flags & GOACC_FLAG_HOST_FALLBACK)) diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90 new file mode 100644 index 00000000000..276a1727b2e --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90 @@ -0,0 +1,197 @@ +! { dg-do run } + +! Test the deviceptr clause with various directives +! and in combination with other directives where +! the deviceptr variable is implied. + +subroutine subr1 (a, b) + implicit none + integer, parameter :: N = 8 + integer :: a(N) + integer :: b(N) + integer :: i = 0 + + !$acc data deviceptr (a) + + !$acc parallel copy (b) + do i = 1, N + a(i) = i * 2 + b(i) = a(i) + end do + !$acc end parallel + + !$acc end data + +end subroutine + +subroutine subr2 (a, b) + implicit none + integer, parameter :: N = 8 + integer :: a(N) + !$acc declare deviceptr (a) + integer :: b(N) + integer :: i = 0 + + !$acc parallel copy (b) + do i = 1, N + a(i) = i * 4 + b(i) = a(i) + end do + !$acc end parallel + +end subroutine + +subroutine subr3 (a, b) + implicit none + integer, parameter :: N = 8 + integer :: a(N) + !$acc declare deviceptr (a) + integer :: b(N) + integer :: i = 0 + + !$acc kernels copy (b) + do i = 1, N + a(i) = i * 8 + b(i) = a(i) + end do + !$acc end kernels + +end subroutine + +subroutine subr4 (a, b) + implicit none + integer, parameter :: N = 8 + integer :: a(N) + integer :: b(N) + integer :: i = 0 + + !$acc parallel deviceptr (a) copy (b) + do i = 1, N + a(i) = i * 16 + b(i) = a(i) + end do + !$acc end parallel + +end subroutine + +subroutine subr5 (a, b) + implicit none + integer, parameter :: N = 8 + integer :: a(N) + integer :: b(N) + integer :: i = 0 + + !$acc kernels deviceptr (a) copy (b) + do i = 1, N + a(i) = i * 32 + b(i) = a(i) + end do + !$acc end kernels + +end subroutine + +subroutine subr6 (a, b) + implicit none + integer, parameter :: N = 8 + integer :: a(N) + integer :: b(N) + integer :: i = 0 + + !$acc parallel deviceptr (a) copy (b) + do i = 1, N + b(i) = i + end do + !$acc end parallel + +end subroutine + +subroutine subr7 (a, b) + implicit none + integer, parameter :: N = 8 + integer :: a(N) + integer :: b(N) + integer :: i = 0 + + !$acc data deviceptr (a) + + !$acc parallel copy (b) + do i = 1, N + a(i) = i * 2 + b(i) = a(i) + end do + !$acc end parallel + + !$acc parallel copy (b) + do i = 1, N + a(i) = b(i) * 2 + b(i) = a(i) + end do + !$acc end parallel + + !$acc end data + +end subroutine + +program main + use iso_c_binding, only: c_ptr, c_f_pointer + implicit none + type (c_ptr) :: cp + integer, parameter :: N = 8 + integer, pointer :: fp(:) + integer :: i = 0 + integer :: b(N) + + interface + function acc_malloc (s) bind (C) + use iso_c_binding, only: c_ptr, c_size_t + integer (c_size_t), value :: s + type (c_ptr) :: acc_malloc + end function + end interface + + cp = acc_malloc (N * sizeof (fp(N))) + call c_f_pointer (cp, fp, [N]) + + call subr1 (fp, b) + + do i = 1, N + if (b(i) .ne. i * 2) call abort + end do + + call subr2 (fp, b) + + do i = 1, N + if (b(i) .ne. i * 4) call abort + end do + + call subr3 (fp, b) + + do i = 1, N + if (b(i) .ne. i * 8) call abort + end do + + call subr4 (fp, b) + + do i = 1, N + if (b(i) .ne. i * 16) call abort + end do + + call subr5 (fp, b) + + do i = 1, N + if (b(i) .ne. i * 32) call abort + end do + + call subr6 (fp, b) + + do i = 1, N + if (b(i) .ne. i) call abort + end do + + call subr7 (fp, b) + + do i = 1, N + if (b(i) .ne. i * 4) call abort + end do + +end program main