https://gcc.gnu.org/g:987bb4c25b4076eb54f043644bdf9988378be90d
commit 987bb4c25b4076eb54f043644bdf9988378be90d
Author: Julian Brown <jul...@codesourcery.com>
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 <ce...@codesourcery.com>
James Norris <jnor...@codesourcery.com>
gcc/fortran/
* openmp.cc (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.cc (gfc_omp_finish_clause): Don't create pointer data
mappings for deviceptr clauses.
(gfc_trans_omp_clauses): Likewise.
gcc/
* gimplify.cc (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/loop-2-kernels-tile.f95: Update.
* gfortran.dg/goacc/loop-2-parallel-tile.f95: Update.
* 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/deviceptr-1.f90: New test.
Diff:
---
gcc/ChangeLog.omp | 10 ++
gcc/fortran/ChangeLog.omp | 6 +
gcc/fortran/openmp.cc | 5 +-
gcc/fortran/trans-openmp.cc | 9 +
gcc/gimplify.cc | 12 +-
gcc/testsuite/ChangeLog.omp | 10 ++
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, 287 insertions(+), 36 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 4b389ee8046..4f4d1dbfad4 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,13 @@
+2018-06-29 Cesar Philippidis <ce...@codesourcery.com>
+ James Norris <jnor...@codesourcery.com>
+
+ * gimplify.cc (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.
+
2023-04-18 Kwok Cheung Yeung <k...@codesourcery.com>
* gimplify.cc (omp_group_base): Handle GOMP_MAP_NONCONTIG_ARRAY_*
diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index b742e1bfe35..379b2aee149 100644
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,9 @@
+2018-06-29 Cesar Philippidis <ce...@codesourcery.com>
+ James Norris <jnor...@codesourcery.com>
+
+ * openmp.cc (resolve_positive_int_expr): Promote the warning to an
+ error.
+
2020-04-19 Chung-Lin Tang <clt...@codesourcery.com>
PR other/76739
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 315ec68f259..2dbe89057b1 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -7158,9 +7158,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 ((flag_openmp || flag_openmp_simd) ? OPT_Wopenmp : 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.cc b/gcc/fortran/trans-openmp.cc
index f867e2240bf..55274e2ff49 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -1553,6 +1553,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)))
@@ -3304,6 +3307,12 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
OMP_CLAUSE_SIZE (node2) = size_int (0);
goto finalize_map_clause;
}
+ else if (POINTER_TYPE_P (type)
+ && n->u.map.op == OMP_MAP_FORCE_DEVICEPTR)
+ {
+ OMP_CLAUSE_DECL (node) = decl;
+ goto finalize_map_clause;
+ }
else if (POINTER_TYPE_P (type)
&& (gfc_omp_privatize_by_reference (decl)
|| GFC_DECL_GET_SCALAR_POINTER (decl)
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 51cad72bf90..12be6d0a6e8 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -151,6 +151,9 @@ enum gimplify_omp_var_data
/* Flag for GOVD_FIRSTPRIVATE: OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT. */
GOVD_FIRSTPRIVATE_IMPLICIT = 0x4000000,
+ /* 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)
@@ -8334,6 +8337,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree
decl, bool in_code)
error ("variable %qE declared in enclosing "
"%<host_data%> 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;
@@ -12689,6 +12693,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_PRESENT_TO
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_PRESENT_TOFROM)
flags |= GOVD_MAP_ALWAYS_TO;
+ else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR)
+ flags |= GOVD_DEVICEPTR;
goto do_add;
@@ -13640,7 +13646,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;
@@ -13666,6 +13673,9 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void
*data)
case GOVD_MAP_FORCE_PRESENT | GOVD_MAP_ALLOC_ONLY:
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 64bb0cb2e5c..75d810faac5 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,13 @@
+2018-06-29 Cesar Philippidis <ce...@codesourcery.com>
+ James Norris <jnor...@codesourcery.com>
+
+ * c-c++-common/goacc/deviceptr-4.c: Update.
+ * gfortran.dg/goacc/loop-2-kernels-tile.f95: Update.
+ * gfortran.dg/goacc/loop-2-parallel-tile.f95: Update.
+ * gfortran.dg/goacc/sie.f95: Update.
+ * gfortran.dg/goacc/tile-1.f90: Update.
+ * gfortran.dg/gomp/pr77516.f90: Update.
+
2020-04-19 Chung-Lin Tang <clt...@codesourcery.com>
PR other/76739
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 5982d5d229f..f393cf29dd4 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 expression after
'num_workers\\('" }
@@ -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 expression after
'num_workers\\('" }
@@ -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 expression after
'vector_length\\('" }
@@ -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 expression after
'vector_length\\('" }
@@ -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 2769eb88dad..fc784d92160 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,11 @@
+2018-06-29 Cesar Philippidis <ce...@codesourcery.com>
+ James Norris <jnor...@codesourcery.com>
+
+ * 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 <jul...@codesourcery.com>
* 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 2c10f07e468..f7222b144e0 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -624,6 +624,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
