public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-11] Add OpenACC Fortran support for deviceptr and variable in common blocks
@ 2021-05-13 16:09 Kwok Yeung
0 siblings, 0 replies; only message in thread
From: Kwok Yeung @ 2021-05-13 16:09 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:a14b3f29681da1d2465e15f98b8cf8d5c64a2c3c
commit a14b3f29681da1d2465e15f98b8cf8d5c64a2c3c
Author: Julian Brown <julian@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 <cesar@codesourcery.com>
James Norris <jnorris@codesourcery.com>
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 <cesar@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+
+ * 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 <thomas@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
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 <cesar@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+
+ * openmp.c (resolve_positive_int_expr): Promote the warning to an
+ error.
+
2018-10-02 Thomas Schwinge <thomas@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
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 "
"%<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;
@@ -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 <cesar@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+
+ * 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 <thomas@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
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 <cesar@codesourcery.com>
+ James Norris <jnorris@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 <julian@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 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
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2021-05-13 16:09 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-05-13 16:09 [gcc/devel/omp/gcc-11] Add OpenACC Fortran support for deviceptr and variable in common blocks Kwok Yeung
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).