* [patch] Add OpenACC Fortran support for deviceptr and variable in common blocks
@ 2018-06-29 19:34 Cesar Philippidis
2018-12-04 13:30 ` Jakub Jelinek
0 siblings, 1 reply; 4+ messages in thread
From: Cesar Philippidis @ 2018-06-29 19:34 UTC (permalink / raw)
To: Fortran List, gcc-patches, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 1583 bytes --]
The attached patch adds support Fortran support for OpenACC deviceptr
and the use of common block variables in data clauses (both implicit and
explicit). This patch also relaxes the Fortran parser to not error
certain types of integral expressions and assumed-sized arrays.
With respect to those errors, I removed them because a lot of working
applications do not explicitly use type attributes (like contiguous).
Perhaps it would be better to reduce them to a warning. Any thoughts on
that? My argument for their removal is that, while the standard states
that, say, arrays must be contiguous or bad things will happen, it does
not necessary mandate that the compiler enforces it. I.e., the intent is
to set the user's expectation that things will go bad if garbage input
is fed to the accelerator. If necessary, I can push back on the OpenACC
standards committee on these issue, but don't expect a quick resolution.
In hindsight, I probably should have kept the error relaxation patches
separate. This patch includes the following patches from og8:
* (dd8b75a) [OpenACC] Update deviceptr handling
* (634727d) [OpenACC] Handle Fortran deviceptr clause
* (d50862a) [Fortran] Remove pointer check in check_array_not_assumed
* (0793cef) [OpenACC] add support for fortran common blocks
* (bdc1acc) [Fortran] update gfortran's tile clause error handling
* (5dc4968) Fix PR72715 "ICE in gfc_trans_omp_do, at
fortran/trans-openmp.c:3164"
Is this patch OK for trunk? It bootstrapped / regression tested cleanly
for x86_64 with nvptx offloading.
Thanks,
Cesar
[-- Attachment #2: 0001-deviceptr-patch-56.patch --]
[-- Type: text/x-patch, Size: 45516 bytes --]
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/loop-5.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.
From 09c1aa87d9a7db2e08384bb47c80b4a61d218a99 Mon Sep 17 00:00:00 2001
From: Cesar Philippidis <cesar@codesourcery.com>
Date: Mon, 25 Jun 2018 13:10:13 -0700
Subject: [PATCH] fortran deviceptr
dd8b75 [OpenACC] Update deviceptr handling
634727 [OpenACC] Handle Fortran deviceptr clause
0793ce [OpenACC] add support for fortran common blocks
bdc1ac [Fortran] update gfortran's tile clause error handling
d50862 [Fortran] Remove pointer check in check_array_not_assumed
5dc496 Fix PR72715 "ICE in gfc_trans_omp_do, at fortran/trans-openmp.c:3164"
---
gcc/fortran/openmp.c | 57 ++---
gcc/fortran/trans-openmp.c | 9 +
gcc/gimplify.c | 35 +++-
.../c-c++-common/goacc/deviceptr-4.c | 2 +-
.../gfortran.dg/goacc/common-block-1.f90 | 69 ++++++
.../gfortran.dg/goacc/common-block-2.f90 | 49 +++++
gcc/testsuite/gfortran.dg/goacc/loop-2.f95 | 8 +-
gcc/testsuite/gfortran.dg/goacc/loop-3-2.f95 | 4 +-
gcc/testsuite/gfortran.dg/goacc/loop-3.f95 | 4 +-
gcc/testsuite/gfortran.dg/goacc/loop-5.f95 | 12 --
gcc/testsuite/gfortran.dg/goacc/pr72715.f90 | 6 +
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/oacc-parallel.c | 11 +-
.../libgomp.oacc-fortran/common-block-1.f90 | 105 ++++++++++
.../libgomp.oacc-fortran/common-block-2.f90 | 150 +++++++++++++
.../libgomp.oacc-fortran/common-block-3.f90 | 137 ++++++++++++
.../libgomp.oacc-fortran/deviceptr-1.f90 | 197 ++++++++++++++++++
19 files changed, 826 insertions(+), 83 deletions(-)
create mode 100644 gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
create mode 100644 gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
create mode 100644 gcc/testsuite/gfortran.dg/goacc/pr72715.f90
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/common-block-1.f90
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/common-block-3.f90
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 94a7f7eaa50..aec2b26deef 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -914,10 +914,11 @@ omp_inv_mask::omp_inv_mask (const omp_mask &m) : omp_mask (m)
mapping. */
static bool
-gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op)
+gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op,
+ bool common_blocks)
{
gfc_omp_namelist **head = NULL;
- if (gfc_match_omp_variable_list ("", list, false, NULL, &head, true)
+ if (gfc_match_omp_variable_list ("", list, common_blocks, NULL, &head, true)
== MATCH_YES)
{
gfc_omp_namelist *n;
@@ -1039,7 +1040,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
if ((mask & OMP_CLAUSE_COPY)
&& gfc_match ("copy ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_TOFROM))
+ OMP_MAP_TOFROM, openacc))
continue;
if (mask & OMP_CLAUSE_COPYIN)
{
@@ -1047,7 +1048,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
{
if (gfc_match ("copyin ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_TO))
+ OMP_MAP_TO, true))
continue;
}
else if (gfc_match_omp_variable_list ("copyin (",
@@ -1058,7 +1059,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
if ((mask & OMP_CLAUSE_COPYOUT)
&& gfc_match ("copyout ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FROM))
+ OMP_MAP_FROM, true))
continue;
if ((mask & OMP_CLAUSE_COPYPRIVATE)
&& gfc_match_omp_variable_list ("copyprivate (",
@@ -1068,7 +1069,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
if ((mask & OMP_CLAUSE_CREATE)
&& gfc_match ("create ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_ALLOC))
+ OMP_MAP_ALLOC, true))
continue;
break;
case 'd':
@@ -1104,7 +1105,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
if ((mask & OMP_CLAUSE_DELETE)
&& gfc_match ("delete ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_RELEASE))
+ OMP_MAP_RELEASE, true))
continue;
if ((mask & OMP_CLAUSE_DEPEND)
&& gfc_match ("depend ( ") == MATCH_YES)
@@ -1156,12 +1157,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
&& openacc
&& gfc_match ("device ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_TO))
+ OMP_MAP_FORCE_TO, false))
continue;
if ((mask & OMP_CLAUSE_DEVICEPTR)
&& gfc_match ("deviceptr ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_DEVICEPTR))
+ OMP_MAP_FORCE_DEVICEPTR, false))
continue;
if ((mask & OMP_CLAUSE_DEVICE_RESIDENT)
&& gfc_match_omp_variable_list
@@ -1239,7 +1240,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
if ((mask & OMP_CLAUSE_HOST_SELF)
&& gfc_match ("host ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_FROM))
+ OMP_MAP_FORCE_FROM, true))
continue;
break;
case 'i':
@@ -1511,47 +1512,47 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
if ((mask & OMP_CLAUSE_COPY)
&& gfc_match ("pcopy ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_TOFROM))
+ OMP_MAP_TOFROM, true))
continue;
if ((mask & OMP_CLAUSE_COPYIN)
&& gfc_match ("pcopyin ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_TO))
+ OMP_MAP_TO, true))
continue;
if ((mask & OMP_CLAUSE_COPYOUT)
&& gfc_match ("pcopyout ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FROM))
+ OMP_MAP_FROM, true))
continue;
if ((mask & OMP_CLAUSE_CREATE)
&& gfc_match ("pcreate ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_ALLOC))
+ OMP_MAP_ALLOC, true))
continue;
if ((mask & OMP_CLAUSE_PRESENT)
&& gfc_match ("present ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_PRESENT))
+ OMP_MAP_FORCE_PRESENT, false))
continue;
if ((mask & OMP_CLAUSE_COPY)
&& gfc_match ("present_or_copy ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_TOFROM))
+ OMP_MAP_TOFROM, true))
continue;
if ((mask & OMP_CLAUSE_COPYIN)
&& gfc_match ("present_or_copyin ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_TO))
+ OMP_MAP_TO, true))
continue;
if ((mask & OMP_CLAUSE_COPYOUT)
&& gfc_match ("present_or_copyout ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FROM))
+ OMP_MAP_FROM, true))
continue;
if ((mask & OMP_CLAUSE_CREATE)
&& gfc_match ("present_or_create ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_ALLOC))
+ OMP_MAP_ALLOC, true))
continue;
if ((mask & OMP_CLAUSE_PRIORITY)
&& c->priority == NULL
@@ -1774,7 +1775,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
if ((mask & OMP_CLAUSE_HOST_SELF)
&& gfc_match ("self ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_FROM))
+ OMP_MAP_FORCE_FROM, true))
continue;
if ((mask & OMP_CLAUSE_SEQ)
&& !c->seq
@@ -3718,8 +3719,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
@@ -3777,10 +3778,6 @@ check_array_not_assumed (gfc_symbol *sym, locus loc, const char *name)
if (sym->as && sym->as->type == AS_ASSUMED_RANK)
gfc_error ("Assumed rank array %qs in %s clause at %L",
sym->name, name, &loc);
- if (sym->as && sym->as->type == AS_DEFERRED && sym->attr.pointer
- && !sym->attr.contiguous)
- gfc_error ("Noncontiguous deferred shape array %qs in %s clause at %L",
- sym->name, name, &loc);
}
static void
@@ -5751,7 +5748,13 @@ resolve_oacc_nested_loops (gfc_code *code, gfc_code* do_code, int collapse,
"at %L", &do_code->loc);
break;
}
- gcc_assert (do_code->op == EXEC_DO || do_code->op == EXEC_DO_CONCURRENT);
+ if (do_code->op == EXEC_DO_CONCURRENT)
+ {
+ gfc_error ("!$ACC LOOP cannot be a DO CONCURRENT loop at %L",
+ &do_code->loc);
+ break;
+ }
+ gcc_assert (do_code->op == EXEC_DO);
if (do_code->ext.iterator->var->ts.type != BT_INTEGER)
gfc_error ("!$ACC LOOP iteration variable must be of type integer at %L",
&do_code->loc);
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index f038f4c5bf8..ca31c88569d 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -1060,6 +1060,8 @@ gfc_omp_finish_clause (tree c, gimple_seq *pre_p)
}
tree c2 = NULL_TREE, c3 = NULL_TREE, c4 = NULL_TREE;
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR)
+ return;
if (POINTER_TYPE_P (TREE_TYPE (decl)))
{
if (!gfc_omp_privatize_by_reference (decl)
@@ -2111,6 +2113,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL)
{
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)
|| GFC_DECL_GET_SCALAR_ALLOCATABLE (decl)
@@ -2282,6 +2290,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
ptr2 = fold_convert (sizetype, ptr2);
OMP_CLAUSE_SIZE (node3)
= fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2);
+ finalize_map_clause:;
}
switch (n->u.map_op)
{
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 48ac92e2b16..346b32031ad 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -105,6 +105,9 @@ enum gimplify_omp_var_data
/* Flag for GOVD_MAP: must be present already. */
GOVD_MAP_FORCE_PRESENT = 524288,
+ /* Flag for OpenACC deviceptrs. */
+ GOVD_DEVICEPTR = (1<<21),
+
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
| GOVD_LOCAL)
@@ -7039,15 +7042,20 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
{
const char *rkind;
bool on_device = false;
+ bool is_private = false;
bool declared = is_oacc_declared (decl);
tree type = TREE_TYPE (decl);
if (lang_hooks.decls.omp_privatize_by_reference (decl))
type = TREE_TYPE (type);
+ if (RECORD_OR_UNION_TYPE_P (type))
+ is_private = lang_hooks.decls.omp_disregard_value_expr (decl, false);
+
if ((ctx->region_type & (ORT_ACC_PARALLEL | ORT_ACC_KERNELS)) != 0
&& is_global_var (decl)
- && device_resident_p (decl))
+ && device_resident_p (decl)
+ && !is_private)
{
on_device = true;
flags |= GOVD_MAP_TO_ONLY;
@@ -7058,7 +7066,9 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
case ORT_ACC_KERNELS:
rkind = "kernels";
- if (AGGREGATE_TYPE_P (type))
+ if (is_private)
+ flags |= GOVD_MAP;
+ else if (AGGREGATE_TYPE_P (type))
{
/* Aggregates default to 'present_or_copy', or 'present'. */
if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
@@ -7075,7 +7085,9 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
case ORT_ACC_PARALLEL:
rkind = "parallel";
- if (on_device || declared)
+ if (is_private)
+ flags |= GOVD_FIRSTPRIVATE;
+ else if (on_device || declared)
flags |= GOVD_MAP;
else if (AGGREGATE_TYPE_P (type))
{
@@ -7141,7 +7153,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
{
tree value = get_base_address (DECL_VALUE_EXPR (decl));
- if (value && DECL_P (value) && DECL_THREAD_LOCAL_P (value))
+ if (!(ctx->region_type & ORT_ACC)
+ && value && DECL_P (value) && DECL_THREAD_LOCAL_P (value))
return omp_notice_threadprivate_variable (ctx, decl, value);
}
@@ -7173,7 +7186,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
if ((ctx->region_type & ORT_TARGET) != 0)
{
- ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
+ shared = !(ctx->region_type & ORT_ACC);
+ ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
if (n == NULL)
{
unsigned nflags = flags;
@@ -7232,6 +7246,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;
@@ -7326,6 +7341,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
}
shared = ((flags | n->value) & GOVD_SHARED) != 0;
+ if (ctx->region_type & ORT_ACC)
+ shared = false;
ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
/* If nothing changed, there's nothing left to do. */
@@ -8213,6 +8230,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;
goto do_add;
case OMP_CLAUSE_DEPEND:
@@ -8828,7 +8847,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
/* Not all combinations of these GOVD_MAP flags are actually valid. */
switch (flags & (GOVD_MAP_TO_ONLY
| GOVD_MAP_FORCE
- | GOVD_MAP_FORCE_PRESENT))
+ | GOVD_MAP_FORCE_PRESENT
+ | GOVD_DEVICEPTR))
{
case 0:
kind = GOMP_MAP_TOFROM;
@@ -8845,6 +8865,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/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/common-block-1.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
new file mode 100644
index 00000000000..c9de125a2f1
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
@@ -0,0 +1,69 @@
+! Test data clauses involving common blocks and common block data.
+! Specifically, validate early matching errors.
+
+subroutine subtest
+ implicit none
+ integer, parameter :: n = 10
+ integer a(n), b(n), c, d(n), e
+ real*4 x(n), y(n), z, w(n), v
+ common /blockA/ a, c, x
+ common /blockB/ b, y, z
+ !$acc declare link(/blockA/, /blockB/, e, v)
+end subroutine subtest
+
+program test
+ implicit none
+ integer, parameter :: n = 10
+ integer a(n), b(n), c, d(n), e
+ real*4 x(n), y(n), z, w(n), v
+ common /blockA/ a, c, x
+ common /blockB/ b, y, z
+ !$acc declare link(/blockA/, /blockB/, e, v)
+
+ !$acc data copy(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data copyin(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data copyout(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data create(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data copyout(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data pcopy(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data pcopyin(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data pcopyout(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data pcreate(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data pcopyout(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data present(/blockA/, /blockB/, e, v) ! { dg-error "" }
+ !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" }
+
+ !$acc parallel private(/blockA/, /blockB/, e, v)
+ !$acc end parallel
+
+ !$acc parallel firstprivate(/blockA/, /blockB/, e, v)
+ !$acc end parallel
+
+ !$acc exit data delete(/blockA/, /blockB/, e, v)
+
+ !$acc data deviceptr(/blockA/, /blockB/, e, v) ! { dg-error "Syntax error" }
+ !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" }
+
+ !$acc data deviceptr(/blockA/, /blockB/, e, v) ! { dg-error "Syntax error" }
+ !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" }
+end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
new file mode 100644
index 00000000000..b83638918a3
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
@@ -0,0 +1,49 @@
+! Test data clauses involving common blocks and common block data.
+! Specifically, resolver errors such as duplicate data clauses.
+
+program test
+ implicit none
+ integer, parameter :: n = 10
+ integer a(n), b(n), c, d(n), e
+ real*4 x(n), y(n), z, w(n), v
+ common /blockA/ a, c, x
+ common /blockB/ b, y, z
+
+ !$acc data copy(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data copyin(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data copyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data create(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data copyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data pcopy(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data pcopyin(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data pcopyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data pcreate(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data pcopyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc parallel private(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end parallel
+
+ !$acc parallel firstprivate(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end parallel
+
+ !$acc exit data delete(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-2.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-2.f95
index 0c902b22410..d4c62732331 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-2.f95
@@ -143,7 +143,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" }
@@ -307,7 +307,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" }
@@ -460,7 +460,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" }
@@ -612,7 +612,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/loop-3-2.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-3-2.f95
index 9be74a85919..c091084a4f5 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-3-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-3-2.f95
@@ -27,9 +27,9 @@ subroutine test1
!$acc end parallel
!$acc end loop ! { dg-error "Unexpected" }
- ! OpenACC supports Fortran 2008 do concurrent statement
+ ! OpenACC does not support Fortran 2008 do concurrent statement
!$acc loop
- do concurrent (i = 1:5)
+ do concurrent (i = 1:5) ! { dg-error "ACC LOOP cannot be a DO CONCURRENT loop" }
end do
!$acc loop
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-3.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-3.f95
index 30930f404f3..ed3e8d50a76 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-3.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-3.f95
@@ -24,9 +24,9 @@ subroutine test1
!$acc end parallel
!$acc end loop ! { dg-error "Unexpected" }
- ! OpenACC supports Fortran 2008 do concurrent statement
+ ! OpenACC does not support Fortran 2008 do concurrent statement
!$acc loop
- do concurrent (i = 1:5)
+ do concurrent (i = 1:5) ! { dg-error "ACC LOOP cannot be a DO CONCURRENT loop" }
end do
!$acc loop
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-5.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95
index d059cf7f377..fe137d515ee 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-5.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95
@@ -93,9 +93,6 @@ program test
DO j = 1,10
ENDDO
ENDDO
- !$acc loop tile(-1) ! { dg-warning "must be positive" }
- do i = 1,10
- enddo
!$acc loop vector tile(*)
DO i = 1,10
ENDDO
@@ -129,9 +126,6 @@ program test
DO j = 1,10
ENDDO
ENDDO
- !$acc loop tile(-1) ! { dg-warning "must be positive" }
- do i = 1,10
- enddo
!$acc loop vector tile(*)
DO i = 1,10
ENDDO
@@ -242,9 +236,6 @@ program test
DO j = 1,10
ENDDO
ENDDO
- !$acc kernels loop tile(-1) ! { dg-warning "must be positive" }
- do i = 1,10
- enddo
!$acc kernels loop vector tile(*)
DO i = 1,10
ENDDO
@@ -333,9 +324,6 @@ program test
DO j = 1,10
ENDDO
ENDDO
- !$acc parallel loop tile(-1) ! { dg-warning "must be positive" }
- do i = 1,10
- enddo
!$acc parallel loop vector tile(*)
DO i = 1,10
ENDDO
diff --git a/gcc/testsuite/gfortran.dg/goacc/pr72715.f90 b/gcc/testsuite/gfortran.dg/goacc/pr72715.f90
new file mode 100644
index 00000000000..68580f9f7de
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/pr72715.f90
@@ -0,0 +1,6 @@
+program p
+ integer :: i
+ !$acc loop
+ do concurrent (i=1:3) ! { dg-error "ACC LOOP cannot be a DO CONCURRENT loop" }
+ end do
+end program p
diff --git a/gcc/testsuite/gfortran.dg/goacc/sie.f95 b/gcc/testsuite/gfortran.dg/goacc/sie.f95
index abfe28bc533..3abf2c8016f 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 3dbabda0342..17fd32cd284 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/oacc-parallel.c b/libgomp/oacc-parallel.c
index b80ace58590..f9cfc4c3a16 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -231,8 +231,13 @@ GOACC_parallel_keyed (int device, void (*fn) (void *),
devaddrs = gomp_alloca (sizeof (void *) * mapnum);
for (i = 0; i < mapnum; i++)
- devaddrs[i] = (void *) (tgt->list[i].key->tgt->tgt_start
- + tgt->list[i].key->tgt_offset);
+ {
+ if (tgt->list[i].key != NULL)
+ devaddrs[i] = (void *) (tgt->list[i].key->tgt->tgt_start
+ + tgt->list[i].key->tgt_offset);
+ else
+ devaddrs[i] = NULL;
+ }
acc_dev->openacc.exec_func (tgt_fn, mapnum, hostaddrs, devaddrs,
async, dims, tgt);
@@ -299,6 +304,8 @@ GOACC_data_start (int device, size_t mapnum,
struct goacc_thread *thr = goacc_thread ();
struct gomp_device_descr *acc_dev = thr->dev;
+ handle_ftn_pointers (mapnum, hostaddrs, sizes, kinds);
+
/* Host fallback or 'do nothing'. */
if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|| host_fallback)
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/common-block-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/common-block-1.f90
new file mode 100644
index 00000000000..9f402973d3d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/common-block-1.f90
@@ -0,0 +1,105 @@
+! Test data located inside common blocks. This test does not execrise
+! ACC DECLARE.
+
+module const
+ integer, parameter :: n = 100
+end module const
+
+subroutine check
+ use const
+
+ implicit none
+ integer i, x(n), y
+ common /BLOCK/ x, y
+
+ do i = 1, n
+ if (x(i) .ne. y) call abort
+ end do
+end subroutine check
+
+module m
+ use const
+ integer a(n), b
+ common /BLOCK/ a, b
+
+contains
+ subroutine mod_implicit_incr
+ implicit none
+ integer i
+
+ !$acc parallel loop
+ do i = 1, n
+ a(i) = b
+ end do
+ !$acc end parallel loop
+
+ call check
+ end subroutine mod_implicit_incr
+
+ subroutine mod_explicit_incr
+ implicit none
+ integer i
+
+ !$acc parallel loop copy(a(1:n)) copyin(b)
+ do i = 1, n
+ a(i) = b
+ end do
+ !$acc end parallel loop
+
+ call check
+ end subroutine mod_explicit_incr
+end module m
+
+subroutine sub_implicit_incr
+ use const
+
+ implicit none
+ integer i, x(n), y
+ common /BLOCK/ x, y
+
+ !$acc parallel loop
+ do i = 1, n
+ x(i) = y
+ end do
+ !$acc end parallel loop
+
+ call check
+end subroutine sub_implicit_incr
+
+subroutine sub_explicit_incr
+ use const
+
+ implicit none
+ integer i, x(n), y
+ common /BLOCK/ x, y
+
+ !$acc parallel loop copy(x(1:n)) copyin(y)
+ do i = 1, n
+ x(i) = y
+ end do
+ !$acc end parallel loop
+
+ call check
+end subroutine sub_explicit_incr
+
+program main
+ use m
+
+ implicit none
+
+ a(:) = -1
+ b = 5
+ call mod_implicit_incr
+
+ a(:) = -2
+ b = 6
+ call mod_explicit_incr
+
+ a(:) = -3
+ b = 7
+ call sub_implicit_incr
+
+ a(:) = -4
+ b = 8
+ call sub_explicit_incr
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90
new file mode 100644
index 00000000000..bf17fc586b9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90
@@ -0,0 +1,150 @@
+! Test data located inside common blocks. This test does not execrise
+! ACC DECLARE. All data clauses are explicit.
+
+module consts
+ integer, parameter :: n = 100
+end module consts
+
+subroutine validate
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 x(n), y(n), z
+ common /BLOCK/ x, y, z, j
+
+ do i = 1, n
+ if (abs(x(i) - i - z) .ge. 0.0001) call abort
+ end do
+end subroutine validate
+
+subroutine incr
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 x(n), y(n), z
+ common /BLOCK/ x, y, z, j
+
+ !$acc parallel loop pcopy(/BLOCK/)
+ do i = 1, n
+ x(i) = x(i) + z
+ end do
+ !$acc end parallel loop
+end subroutine incr
+
+program main
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 a(n), b(n), c
+ common /BLOCK/ a, b, c, j
+
+ ! Test copyout, pcopy, device
+
+ !$acc data copyout(a, c)
+
+ c = 1.0
+
+ !$acc update device(c)
+
+ !$acc parallel loop pcopy(a)
+ do i = 1, n
+ a(i) = i
+ end do
+ !$acc end parallel loop
+
+ call incr
+ call incr
+ call incr
+ !$acc end data
+
+ c = 3.0
+ call validate
+
+ ! Test pcopy without copyout
+
+ c = 2.0
+ call incr
+ c = 5.0
+ call validate
+
+ ! Test create, delete, host, copyout, copyin
+
+ !$acc enter data create(b)
+
+ !$acc parallel loop pcopy(b)
+ do i = 1, n
+ b(i) = i
+ end do
+ !$acc end parallel loop
+
+ !$acc update host (b)
+
+ !$acc parallel loop pcopy(b) copyout(a) copyin(c)
+ do i = 1, n
+ a(i) = b(i) + c
+ end do
+ !$acc end parallel loop
+
+ !$acc exit data delete(b)
+
+ call validate
+
+ a(:) = b(:)
+ c = 0.0
+ call validate
+
+ ! Test copy
+
+ c = 1.0
+ !$acc parallel loop copy(/BLOCK/)
+ do i = 1, n
+ a(i) = b(i) + c
+ end do
+ !$acc end parallel loop
+
+ call validate
+
+ ! Test pcopyin, pcopyout FIXME
+
+ c = 2.0
+ !$acc data copyin(b, c) copyout(a)
+
+ !$acc parallel loop pcopyin(b, c) pcopyout(a)
+ do i = 1, n
+ a(i) = b(i) + c
+ end do
+ !$acc end parallel loop
+
+ !$acc end data
+
+ call validate
+
+ ! Test reduction, private
+
+ j = 0
+
+ !$acc parallel private(i) copy(j)
+ !$acc loop reduction(+:j)
+ do i = 1, n
+ j = j + 1
+ end do
+ !$acc end parallel
+
+ if (j .ne. n) call abort
+
+ ! Test firstprivate, copy
+
+ a(:) = 0
+ c = j
+
+ !$acc parallel loop firstprivate(c) copyout(a)
+ do i = 1, n
+ a(i) = i + c
+ end do
+ !$acc end parallel loop
+
+ call validate
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/common-block-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/common-block-3.f90
new file mode 100644
index 00000000000..134e2d1cf29
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/common-block-3.f90
@@ -0,0 +1,137 @@
+! Test data located inside common blocks. This test does not execrise
+! ACC DECLARE. Most of the data clauses are implicit.
+
+module consts
+ integer, parameter :: n = 100
+end module consts
+
+subroutine validate
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 x(n), y(n), z
+ common /BLOCK/ x, y, z, j
+
+ do i = 1, n
+ if (abs(x(i) - i - z) .ge. 0.0001) call abort
+ end do
+end subroutine validate
+
+subroutine incr_parallel
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 x(n), y(n), z
+ common /BLOCK/ x, y, z, j
+
+ !$acc parallel loop
+ do i = 1, n
+ x(i) = x(i) + z
+ end do
+ !$acc end parallel loop
+end subroutine incr_parallel
+
+subroutine incr_kernels
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 x(n), y(n), z
+ common /BLOCK/ x, y, z, j
+
+ !$acc kernels
+ do i = 1, n
+ x(i) = x(i) + z
+ end do
+ !$acc end kernels
+end subroutine incr_kernels
+
+program main
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 a(n), b(n), c
+ common /BLOCK/ a, b, c, j
+
+ !$acc data copyout(a, c)
+
+ c = 1.0
+
+ !$acc update device(c)
+
+ !$acc parallel loop
+ do i = 1, n
+ a(i) = i
+ end do
+ !$acc end parallel loop
+
+ call incr_parallel
+ call incr_parallel
+ call incr_parallel
+ !$acc end data
+
+ c = 3.0
+ call validate
+
+ ! Test pcopy without copyout
+
+ c = 2.0
+ call incr_kernels
+ c = 5.0
+ call validate
+
+ !$acc kernels
+ do i = 1, n
+ b(i) = i
+ end do
+ !$acc end kernels
+
+ !$acc parallel loop
+ do i = 1, n
+ a(i) = b(i) + c
+ end do
+ !$acc end parallel loop
+
+ call validate
+
+ a(:) = b(:)
+ c = 0.0
+ call validate
+
+ ! Test copy
+
+ c = 1.0
+ !$acc parallel loop
+ do i = 1, n
+ a(i) = b(i) + c
+ end do
+ !$acc end parallel loop
+
+ call validate
+
+ c = 2.0
+ !$acc data copyin(b, c) copyout(a)
+
+ !$acc kernels
+ do i = 1, n
+ a(i) = b(i) + c
+ end do
+ !$acc end kernels
+
+ !$acc end data
+
+ call validate
+
+ j = 0
+
+ !$acc parallel loop reduction(+:j)
+ do i = 1, n
+ j = j + 1
+ end do
+ !$acc end parallel loop
+
+ if (j .ne. n) call abort
+end program main
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
--
2.17.1
^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [patch] Add OpenACC Fortran support for deviceptr and variable in common blocks
2018-06-29 19:34 [patch] Add OpenACC Fortran support for deviceptr and variable in common blocks Cesar Philippidis
@ 2018-12-04 13:30 ` Jakub Jelinek
2018-12-04 13:50 ` Chung-Lin Tang
0 siblings, 1 reply; 4+ messages in thread
From: Jakub Jelinek @ 2018-12-04 13:30 UTC (permalink / raw)
To: Cesar Philippidis, Thomas Schwinge; +Cc: Fortran List, gcc-patches
On Fri, Jun 29, 2018 at 12:04:58PM -0700, Cesar Philippidis wrote:
> 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/loop-5.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.
> --- a/gcc/fortran/openmp.c
> +++ b/gcc/fortran/openmp.c
> @@ -914,10 +914,11 @@ omp_inv_mask::omp_inv_mask (const omp_mask &m) : omp_mask (m)
> mapping. */
>
> static bool
> -gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op)
> +gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op,
> + bool common_blocks)
> {
> gfc_omp_namelist **head = NULL;
> - if (gfc_match_omp_variable_list ("", list, false, NULL, &head, true)
> + if (gfc_match_omp_variable_list ("", list, common_blocks, NULL, &head, true)
> == MATCH_YES)
> {
> gfc_omp_namelist *n;
> @@ -1039,7 +1040,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
> if ((mask & OMP_CLAUSE_COPY)
> && gfc_match ("copy ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_TOFROM))
> + OMP_MAP_TOFROM, openacc))
Why? OpenMP doesn't have a copy clause, so I'd expect true here.
> continue;
> if (mask & OMP_CLAUSE_COPYIN)
> {
> @@ -1047,7 +1048,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
> {
> if (gfc_match ("copyin ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_TO))
> + OMP_MAP_TO, true))
> continue;
OpenMP does have a copyin clause, but it is handled below, so this one is ok.
> @@ -1156,12 +1157,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
> && openacc
> && gfc_match ("device ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_FORCE_TO))
> + OMP_MAP_FORCE_TO, false))
> continue;
> if ((mask & OMP_CLAUSE_DEVICEPTR)
> && gfc_match ("deviceptr ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_FORCE_DEVICEPTR))
> + OMP_MAP_FORCE_DEVICEPTR, false))
> continue;
Not sure about these, your call. deviceptr is OpenACC specific clause,
device is in both, but means something different in OpenMP. In any case,
haven't looked if OpenACC allows common blocks in these clauses.
> @@ -3718,8 +3719,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);
> }
This affects OpenMP too and makes it inconsistent with C/C++. Why?
If you need it for OpenACC clauses, then we need two routines.
> static void
> @@ -3777,10 +3778,6 @@ check_array_not_assumed (gfc_symbol *sym, locus loc, const char *name)
> if (sym->as && sym->as->type == AS_ASSUMED_RANK)
> gfc_error ("Assumed rank array %qs in %s clause at %L",
> sym->name, name, &loc);
> - if (sym->as && sym->as->type == AS_DEFERRED && sym->attr.pointer
> - && !sym->attr.contiguous)
> - gfc_error ("Noncontiguous deferred shape array %qs in %s clause at %L",
> - sym->name, name, &loc);
> }
Perhaps it would help to mention oacc in the name of this routine to make
it clearer that it is OpenACC only. The change is ok if that is what
OpenACC now allows.
> --- a/gcc/gimplify.c
> +++ b/gcc/gimplify.c
> @@ -105,6 +105,9 @@ enum gimplify_omp_var_data
> /* Flag for GOVD_MAP: must be present already. */
> GOVD_MAP_FORCE_PRESENT = 524288,
>
> + /* Flag for OpenACC deviceptrs. */
> + GOVD_DEVICEPTR = (1<<21),
Please use the same style of constants as in the rest (and, you need
to adjust anyway for current trunk).
Jakub
^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [patch] Add OpenACC Fortran support for deviceptr and variable in common blocks
2018-12-04 13:30 ` Jakub Jelinek
@ 2018-12-04 13:50 ` Chung-Lin Tang
2019-01-02 14:38 ` Chung-Lin Tang
0 siblings, 1 reply; 4+ messages in thread
From: Chung-Lin Tang @ 2018-12-04 13:50 UTC (permalink / raw)
To: Jakub Jelinek, Cesar Philippidis, Thomas Schwinge
Cc: Fortran List, gcc-patches, Moore, Catherine
On 2018/12/4 9:30 PM, Jakub Jelinek wrote:
>> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
>> - OMP_MAP_TOFROM))
>> + OMP_MAP_TOFROM, openacc))
> Why? OpenMP doesn't have a copy clause, so I'd expect true here.
I guess Cesar was just being specific to OpenACC behavior.
>> continue;
>> if (mask & OMP_CLAUSE_COPYIN)
>> {
>> @@ -1047,7 +1048,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
>> {
>> if (gfc_match ("copyin ( ") == MATCH_YES
>> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
>> - OMP_MAP_TO))
>> + OMP_MAP_TO, true))
>> continue;
> OpenMP does have a copyin clause, but it is handled below, so this one is ok.
>
>> @@ -1156,12 +1157,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
>> && openacc
>> && gfc_match ("device ( ") == MATCH_YES
>> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
>> - OMP_MAP_FORCE_TO))
>> + OMP_MAP_FORCE_TO, false))
>> continue;
>> if ((mask & OMP_CLAUSE_DEVICEPTR)
>> && gfc_match ("deviceptr ( ") == MATCH_YES
>> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
>> - OMP_MAP_FORCE_DEVICEPTR))
>> + OMP_MAP_FORCE_DEVICEPTR, false))
>> continue;
> Not sure about these, your call. deviceptr is OpenACC specific clause,
> device is in both, but means something different in OpenMP. In any case,
> haven't looked if OpenACC allows common blocks in these clauses.
Common block support and the deviceptr changes were originally separate patches.
Also, I think above changes probably could be greatly shortened if a C++ default
argument were used for gfc_match_omp_map_clause().
>> @@ -3718,8 +3719,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);
>> }
> This affects OpenMP too and makes it inconsistent with C/C++. Why?
> If you need it for OpenACC clauses, then we need two routines.
We'll update this.
>> --- a/gcc/gimplify.c
>> +++ b/gcc/gimplify.c
>> @@ -105,6 +105,9 @@ enum gimplify_omp_var_data
>> /* Flag for GOVD_MAP: must be present already. */
>> GOVD_MAP_FORCE_PRESENT = 524288,
>>
>> + /* Flag for OpenACC deviceptrs. */
>> + GOVD_DEVICEPTR = (1<<21),
> Please use the same style of constants as in the rest (and, you need
> to adjust anyway for current trunk).
This will be updated too.
Thanks,
Chung-Lin
^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [patch] Add OpenACC Fortran support for deviceptr and variable in common blocks
2018-12-04 13:50 ` Chung-Lin Tang
@ 2019-01-02 14:38 ` Chung-Lin Tang
0 siblings, 0 replies; 4+ messages in thread
From: Chung-Lin Tang @ 2019-01-02 14:38 UTC (permalink / raw)
To: Jakub Jelinek, Cesar Philippidis, Thomas Schwinge
Cc: Fortran List, gcc-patches, Moore, Catherine
[-- Attachment #1: Type: text/plain, Size: 205 bytes --]
Hi Jakub,
I have updated this patch and rebased with current trunk. The issues you identified
in the last review should all be resolved. Is this now okay for trunk?
Thanks, and Happy New Year~
Chung-Lin
[-- Attachment #2: openacc-fortran-deviceptr-commonblk.v2.patch --]
[-- Type: text/plain, Size: 39083 bytes --]
Index: gcc/fortran/openmp.c
===================================================================
--- gcc/fortran/openmp.c (revision 267513)
+++ gcc/fortran/openmp.c (working copy)
@@ -914,10 +914,11 @@ omp_inv_mask::omp_inv_mask (const omp_mask &m) : o
mapping. */
static bool
-gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op)
+gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op,
+ bool common_blocks = true)
{
gfc_omp_namelist **head = NULL;
- if (gfc_match_omp_variable_list ("", list, false, NULL, &head, true)
+ if (gfc_match_omp_variable_list ("", list, common_blocks, NULL, &head, true)
== MATCH_YES)
{
gfc_omp_namelist *n;
@@ -1156,12 +1157,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const
&& openacc
&& gfc_match ("device ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_TO))
+ OMP_MAP_FORCE_TO, false))
continue;
if ((mask & OMP_CLAUSE_DEVICEPTR)
&& gfc_match ("deviceptr ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_DEVICEPTR))
+ OMP_MAP_FORCE_DEVICEPTR, false))
continue;
if ((mask & OMP_CLAUSE_DEVICE_RESIDENT)
&& gfc_match_omp_variable_list
@@ -1531,7 +1532,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const
if ((mask & OMP_CLAUSE_PRESENT)
&& gfc_match ("present ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_PRESENT))
+ OMP_MAP_FORCE_PRESENT, false))
continue;
if ((mask & OMP_CLAUSE_COPY)
&& gfc_match ("present_or_copy ( ") == MATCH_YES
@@ -3762,7 +3763,7 @@ check_symbol_not_pointer (gfc_symbol *sym, locus l
/* Emits error when symbol represents assumed size/rank array. */
static void
-check_array_not_assumed (gfc_symbol *sym, locus loc, const char *name)
+check_oacc_array_not_assumed (gfc_symbol *sym, locus loc, const char *name)
{
if (sym->as && sym->as->type == AS_ASSUMED_SIZE)
gfc_error ("Assumed size array %qs in %s clause at %L",
@@ -3770,10 +3771,6 @@ static void
if (sym->as && sym->as->type == AS_ASSUMED_RANK)
gfc_error ("Assumed rank array %qs in %s clause at %L",
sym->name, name, &loc);
- if (sym->as && sym->as->type == AS_DEFERRED && sym->attr.pointer
- && !sym->attr.contiguous)
- gfc_error ("Noncontiguous deferred shape array %qs in %s clause at %L",
- sym->name, name, &loc);
}
static void
@@ -3788,7 +3785,7 @@ resolve_oacc_data_clauses (gfc_symbol *sym, locus
gfc_error ("ALLOCATABLE object %qs of polymorphic type "
"in %s clause at %L", sym->name, name, &loc);
check_symbol_not_pointer (sym, loc, name);
- check_array_not_assumed (sym, loc, name);
+ check_oacc_array_not_assumed (sym, loc, name);
}
static void
@@ -3817,7 +3814,7 @@ resolve_oacc_deviceptr_clause (gfc_symbol *sym, lo
if (sym->attr.value)
gfc_error ("VALUE object %qs in %s clause at %L",
sym->name, name, &loc);
- check_array_not_assumed (sym, loc, name);
+ check_oacc_array_not_assumed (sym, loc, name);
}
@@ -4508,7 +4505,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_claus
}
if (code
&& (oacc_is_loop (code) || code->op == EXEC_OACC_PARALLEL))
- check_array_not_assumed (n->sym, n->where, name);
+ check_oacc_array_not_assumed (n->sym, n->where, name);
else if (n->sym->as && n->sym->as->type == AS_ASSUMED_SIZE)
gfc_error ("Assumed size array %qs in %s clause at %L",
n->sym->name, name, &n->where);
@@ -4725,7 +4722,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_claus
/* FALLTHRU */
case OMP_LIST_DEVICE_RESIDENT:
check_symbol_not_pointer (n->sym, n->where, name);
- check_array_not_assumed (n->sym, n->where, name);
+ check_oacc_array_not_assumed (n->sym, n->where, name);
break;
default:
break;
@@ -5760,7 +5757,13 @@ resolve_oacc_nested_loops (gfc_code *code, gfc_cod
"at %L", &do_code->loc);
break;
}
- gcc_assert (do_code->op == EXEC_DO || do_code->op == EXEC_DO_CONCURRENT);
+ if (do_code->op == EXEC_DO_CONCURRENT)
+ {
+ gfc_error ("!$ACC LOOP cannot be a DO CONCURRENT loop at %L",
+ &do_code->loc);
+ break;
+ }
+ gcc_assert (do_code->op == EXEC_DO);
if (do_code->ext.iterator->var->ts.type != BT_INTEGER)
gfc_error ("!$ACC LOOP iteration variable must be of type integer at %L",
&do_code->loc);
@@ -5878,7 +5881,12 @@ resolve_oacc_loop_blocks (gfc_code *code)
}
else
{
- resolve_positive_int_expr (el->expr, "TILE");
+ resolve_scalar_int_expr (el->expr, "TILE");
+ if (el->expr->expr_type == EXPR_CONSTANT
+ && el->expr->ts.type == BT_INTEGER
+ && mpz_sgn (el->expr->value.integer) <= 0)
+ gfc_error ("INTEGER expression of TILE clause at %L must be "
+ "positive", &el->expr->where);
if (el->expr->expr_type != EXPR_CONSTANT)
gfc_error ("TILE requires constant expression at %L",
&code->loc);
@@ -5965,7 +5973,7 @@ gfc_resolve_oacc_declare (gfc_namespace *ns)
}
for (n = oc->clauses->lists[OMP_LIST_DEVICE_RESIDENT]; n; n = n->next)
- check_array_not_assumed (n->sym, oc->loc, "DEVICE_RESIDENT");
+ check_oacc_array_not_assumed (n->sym, oc->loc, "DEVICE_RESIDENT");
}
for (oc = ns->oacc_declare; oc; oc = oc->next)
Index: gcc/fortran/trans-openmp.c
===================================================================
--- gcc/fortran/trans-openmp.c (revision 267513)
+++ gcc/fortran/trans-openmp.c (working copy)
@@ -1066,6 +1066,8 @@ gfc_omp_finish_clause (tree c, gimple_seq *pre_p)
}
tree c2 = NULL_TREE, c3 = NULL_TREE, c4 = NULL_TREE;
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR)
+ return;
if (POINTER_TYPE_P (TREE_TYPE (decl)))
{
if (!gfc_omp_privatize_by_reference (decl)
@@ -2117,6 +2119,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp
if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL)
{
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)
|| GFC_DECL_GET_SCALAR_ALLOCATABLE (decl)
@@ -2288,6 +2296,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp
ptr2 = fold_convert (sizetype, ptr2);
OMP_CLAUSE_SIZE (node3)
= fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2);
+ finalize_map_clause:;
}
switch (n->u.map_op)
{
Index: gcc/gimplify.c
===================================================================
--- gcc/gimplify.c (revision 267513)
+++ gcc/gimplify.c (working copy)
@@ -113,6 +113,9 @@ enum gimplify_omp_var_data
GOVD_NONTEMPORAL = 4194304,
+ /* Flag for OpenACC deviceptrs. */
+ GOVD_DEVICEPTR = 8388608,
+
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
| GOVD_LOCAL)
@@ -7109,6 +7112,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx,
{
const char *rkind;
bool on_device = false;
+ bool is_private = false;
bool declared = is_oacc_declared (decl);
tree type = TREE_TYPE (decl);
@@ -7115,9 +7119,13 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx,
if (lang_hooks.decls.omp_privatize_by_reference (decl))
type = TREE_TYPE (type);
+ if (RECORD_OR_UNION_TYPE_P (type))
+ is_private = lang_hooks.decls.omp_disregard_value_expr (decl, false);
+
if ((ctx->region_type & (ORT_ACC_PARALLEL | ORT_ACC_KERNELS)) != 0
&& is_global_var (decl)
- && device_resident_p (decl))
+ && device_resident_p (decl)
+ && !is_private)
{
on_device = true;
flags |= GOVD_MAP_TO_ONLY;
@@ -7128,7 +7136,9 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx,
case ORT_ACC_KERNELS:
rkind = "kernels";
- if (AGGREGATE_TYPE_P (type))
+ if (is_private)
+ flags |= GOVD_MAP;
+ else if (AGGREGATE_TYPE_P (type))
{
/* Aggregates default to 'present_or_copy', or 'present'. */
if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
@@ -7145,7 +7155,9 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx,
case ORT_ACC_PARALLEL:
rkind = "parallel";
- if (on_device || declared)
+ if (is_private)
+ flags |= GOVD_FIRSTPRIVATE;
+ else if (on_device || declared)
flags |= GOVD_MAP;
else if (AGGREGATE_TYPE_P (type))
{
@@ -7211,7 +7223,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx,
{
tree value = get_base_address (DECL_VALUE_EXPR (decl));
- if (value && DECL_P (value) && DECL_THREAD_LOCAL_P (value))
+ if (!(ctx->region_type & ORT_ACC)
+ && value && DECL_P (value) && DECL_THREAD_LOCAL_P (value))
return omp_notice_threadprivate_variable (ctx, decl, value);
}
@@ -7243,7 +7256,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx,
n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
if ((ctx->region_type & ORT_TARGET) != 0)
{
- ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
+ shared = !(ctx->region_type & ORT_ACC);
+ ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
if (n == NULL)
{
unsigned nflags = flags;
@@ -7316,6 +7330,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx,
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;
@@ -7411,6 +7426,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx,
}
shared = ((flags | n->value) & GOVD_SHARED) != 0;
+ if (ctx->region_type & ORT_ACC)
+ shared = false;
ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
/* If nothing changed, there's nothing left to do. */
@@ -8804,6 +8821,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_se
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;
goto do_add;
case OMP_CLAUSE_DEPEND:
@@ -9534,7 +9553,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n,
| 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;
@@ -9557,6 +9577,9 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n,
case GOVD_MAP_FORCE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
+ case GOVD_DEVICEPTR:
+ kind = GOMP_MAP_FORCE_DEVICEPTR;
+ break;
default:
gcc_unreachable ();
}
Index: gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/deviceptr-4.c (revision 267513)
+++ gcc/testsuite/c-c++-common/goacc/deviceptr-4.c (working copy)
@@ -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" } } */
Index: gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/common-block-1.f90 (nonexistent)
+++ gcc/testsuite/gfortran.dg/goacc/common-block-1.f90 (working copy)
@@ -0,0 +1,69 @@
+! Test data clauses involving common blocks and common block data.
+! Specifically, validate early matching errors.
+
+subroutine subtest
+ implicit none
+ integer, parameter :: n = 10
+ integer a(n), b(n), c, d(n), e
+ real*4 x(n), y(n), z, w(n), v
+ common /blockA/ a, c, x
+ common /blockB/ b, y, z
+ !$acc declare link(/blockA/, /blockB/, e, v)
+end subroutine subtest
+
+program test
+ implicit none
+ integer, parameter :: n = 10
+ integer a(n), b(n), c, d(n), e
+ real*4 x(n), y(n), z, w(n), v
+ common /blockA/ a, c, x
+ common /blockB/ b, y, z
+ !$acc declare link(/blockA/, /blockB/, e, v)
+
+ !$acc data copy(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data copyin(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data copyout(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data create(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data copyout(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data pcopy(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data pcopyin(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data pcopyout(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data pcreate(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data pcopyout(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data present(/blockA/, /blockB/, e, v) ! { dg-error "" }
+ !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" }
+
+ !$acc parallel private(/blockA/, /blockB/, e, v)
+ !$acc end parallel
+
+ !$acc parallel firstprivate(/blockA/, /blockB/, e, v)
+ !$acc end parallel
+
+ !$acc exit data delete(/blockA/, /blockB/, e, v)
+
+ !$acc data deviceptr(/blockA/, /blockB/, e, v) ! { dg-error "Syntax error" }
+ !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" }
+
+ !$acc data deviceptr(/blockA/, /blockB/, e, v) ! { dg-error "Syntax error" }
+ !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" }
+end program test
Index: gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 (nonexistent)
+++ gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 (working copy)
@@ -0,0 +1,49 @@
+! Test data clauses involving common blocks and common block data.
+! Specifically, resolver errors such as duplicate data clauses.
+
+program test
+ implicit none
+ integer, parameter :: n = 10
+ integer a(n), b(n), c, d(n), e
+ real*4 x(n), y(n), z, w(n), v
+ common /blockA/ a, c, x
+ common /blockB/ b, y, z
+
+ !$acc data copy(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data copyin(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data copyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data create(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data copyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data pcopy(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data pcopyin(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data pcopyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data pcreate(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data pcopyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc parallel private(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end parallel
+
+ !$acc parallel firstprivate(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end parallel
+
+ !$acc exit data delete(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+end program test
Index: gcc/testsuite/gfortran.dg/goacc/loop-2-kernels-tile.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/loop-2-kernels-tile.f95 (revision 267513)
+++ gcc/testsuite/gfortran.dg/goacc/loop-2-kernels-tile.f95 (working copy)
@@ -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" }
Index: gcc/testsuite/gfortran.dg/goacc/loop-2-parallel-tile.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/loop-2-parallel-tile.f95 (revision 267513)
+++ gcc/testsuite/gfortran.dg/goacc/loop-2-parallel-tile.f95 (working copy)
@@ -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" }
Index: gcc/testsuite/gfortran.dg/goacc/loop-3-2.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/loop-3-2.f95 (revision 267513)
+++ gcc/testsuite/gfortran.dg/goacc/loop-3-2.f95 (working copy)
@@ -27,9 +27,9 @@ subroutine test1
!$acc end parallel
!$acc end loop ! { dg-error "Unexpected" }
- ! OpenACC supports Fortran 2008 do concurrent statement
+ ! OpenACC does not support Fortran 2008 do concurrent statement
!$acc loop
- do concurrent (i = 1:5)
+ do concurrent (i = 1:5) ! { dg-error "ACC LOOP cannot be a DO CONCURRENT loop" }
end do
!$acc loop
Index: gcc/testsuite/gfortran.dg/goacc/loop-3.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/loop-3.f95 (revision 267513)
+++ gcc/testsuite/gfortran.dg/goacc/loop-3.f95 (working copy)
@@ -24,9 +24,9 @@ subroutine test1
!$acc end parallel
!$acc end loop ! { dg-error "Unexpected" }
- ! OpenACC supports Fortran 2008 do concurrent statement
+ ! OpenACC does not support Fortran 2008 do concurrent statement
!$acc loop
- do concurrent (i = 1:5)
+ do concurrent (i = 1:5) ! { dg-error "ACC LOOP cannot be a DO CONCURRENT loop" }
end do
!$acc loop
Index: gcc/testsuite/gfortran.dg/goacc/loop-5.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/loop-5.f95 (revision 267513)
+++ gcc/testsuite/gfortran.dg/goacc/loop-5.f95 (working copy)
@@ -93,9 +93,6 @@ program test
DO j = 1,10
ENDDO
ENDDO
- !$acc loop tile(-1) ! { dg-warning "must be positive" }
- do i = 1,10
- enddo
!$acc loop vector tile(*)
DO i = 1,10
ENDDO
@@ -129,9 +126,6 @@ program test
DO j = 1,10
ENDDO
ENDDO
- !$acc loop tile(-1) ! { dg-warning "must be positive" }
- do i = 1,10
- enddo
!$acc loop vector tile(*)
DO i = 1,10
ENDDO
@@ -242,9 +236,6 @@ program test
DO j = 1,10
ENDDO
ENDDO
- !$acc kernels loop tile(-1) ! { dg-warning "must be positive" }
- do i = 1,10
- enddo
!$acc kernels loop vector tile(*)
DO i = 1,10
ENDDO
@@ -333,9 +324,6 @@ program test
DO j = 1,10
ENDDO
ENDDO
- !$acc parallel loop tile(-1) ! { dg-warning "must be positive" }
- do i = 1,10
- enddo
!$acc parallel loop vector tile(*)
DO i = 1,10
ENDDO
Index: gcc/testsuite/gfortran.dg/goacc/pr72715.f90
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/pr72715.f90 (nonexistent)
+++ gcc/testsuite/gfortran.dg/goacc/pr72715.f90 (working copy)
@@ -0,0 +1,6 @@
+program p
+ integer :: i
+ !$acc loop
+ do concurrent (i=1:3) ! { dg-error "ACC LOOP cannot be a DO CONCURRENT loop" }
+ end do
+end program p
Index: gcc/testsuite/gfortran.dg/goacc/sie.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/sie.f95 (revision 267513)
+++ gcc/testsuite/gfortran.dg/goacc/sie.f95 (working copy)
@@ -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" }
Index: gcc/testsuite/gfortran.dg/goacc/tile-1.f90
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/tile-1.f90 (revision 267513)
+++ gcc/testsuite/gfortran.dg/goacc/tile-1.f90 (working copy)
@@ -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
Index: gcc/testsuite/gfortran.dg/gomp/pr77516.f90
===================================================================
--- gcc/testsuite/gfortran.dg/gomp/pr77516.f90 (revision 267513)
+++ gcc/testsuite/gfortran.dg/gomp/pr77516.f90 (working copy)
@@ -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
Index: libgomp/oacc-parallel.c
===================================================================
--- libgomp/oacc-parallel.c (revision 267513)
+++ libgomp/oacc-parallel.c (working copy)
@@ -312,6 +312,8 @@ GOACC_data_start (int flags_m, size_t mapnum,
struct goacc_thread *thr = goacc_thread ();
struct gomp_device_descr *acc_dev = thr->dev;
+ 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))
Index: libgomp/testsuite/libgomp.oacc-fortran/common-block-1.f90
===================================================================
--- libgomp/testsuite/libgomp.oacc-fortran/common-block-1.f90 (nonexistent)
+++ libgomp/testsuite/libgomp.oacc-fortran/common-block-1.f90 (working copy)
@@ -0,0 +1,105 @@
+! Test data located inside common blocks. This test does not execrise
+! ACC DECLARE.
+
+module const
+ integer, parameter :: n = 100
+end module const
+
+subroutine check
+ use const
+
+ implicit none
+ integer i, x(n), y
+ common /BLOCK/ x, y
+
+ do i = 1, n
+ if (x(i) .ne. y) call abort
+ end do
+end subroutine check
+
+module m
+ use const
+ integer a(n), b
+ common /BLOCK/ a, b
+
+contains
+ subroutine mod_implicit_incr
+ implicit none
+ integer i
+
+ !$acc parallel loop
+ do i = 1, n
+ a(i) = b
+ end do
+ !$acc end parallel loop
+
+ call check
+ end subroutine mod_implicit_incr
+
+ subroutine mod_explicit_incr
+ implicit none
+ integer i
+
+ !$acc parallel loop copy(a(1:n)) copyin(b)
+ do i = 1, n
+ a(i) = b
+ end do
+ !$acc end parallel loop
+
+ call check
+ end subroutine mod_explicit_incr
+end module m
+
+subroutine sub_implicit_incr
+ use const
+
+ implicit none
+ integer i, x(n), y
+ common /BLOCK/ x, y
+
+ !$acc parallel loop
+ do i = 1, n
+ x(i) = y
+ end do
+ !$acc end parallel loop
+
+ call check
+end subroutine sub_implicit_incr
+
+subroutine sub_explicit_incr
+ use const
+
+ implicit none
+ integer i, x(n), y
+ common /BLOCK/ x, y
+
+ !$acc parallel loop copy(x(1:n)) copyin(y)
+ do i = 1, n
+ x(i) = y
+ end do
+ !$acc end parallel loop
+
+ call check
+end subroutine sub_explicit_incr
+
+program main
+ use m
+
+ implicit none
+
+ a(:) = -1
+ b = 5
+ call mod_implicit_incr
+
+ a(:) = -2
+ b = 6
+ call mod_explicit_incr
+
+ a(:) = -3
+ b = 7
+ call sub_implicit_incr
+
+ a(:) = -4
+ b = 8
+ call sub_explicit_incr
+end program main
Index: libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90
===================================================================
--- libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90 (nonexistent)
+++ libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90 (working copy)
@@ -0,0 +1,150 @@
+! Test data located inside common blocks. This test does not execrise
+! ACC DECLARE. All data clauses are explicit.
+
+module consts
+ integer, parameter :: n = 100
+end module consts
+
+subroutine validate
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 x(n), y(n), z
+ common /BLOCK/ x, y, z, j
+
+ do i = 1, n
+ if (abs(x(i) - i - z) .ge. 0.0001) call abort
+ end do
+end subroutine validate
+
+subroutine incr
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 x(n), y(n), z
+ common /BLOCK/ x, y, z, j
+
+ !$acc parallel loop pcopy(/BLOCK/)
+ do i = 1, n
+ x(i) = x(i) + z
+ end do
+ !$acc end parallel loop
+end subroutine incr
+
+program main
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 a(n), b(n), c
+ common /BLOCK/ a, b, c, j
+
+ ! Test copyout, pcopy, device
+
+ !$acc data copyout(a, c)
+
+ c = 1.0
+
+ !$acc update device(c)
+
+ !$acc parallel loop pcopy(a)
+ do i = 1, n
+ a(i) = i
+ end do
+ !$acc end parallel loop
+
+ call incr
+ call incr
+ call incr
+ !$acc end data
+
+ c = 3.0
+ call validate
+
+ ! Test pcopy without copyout
+
+ c = 2.0
+ call incr
+ c = 5.0
+ call validate
+
+ ! Test create, delete, host, copyout, copyin
+
+ !$acc enter data create(b)
+
+ !$acc parallel loop pcopy(b)
+ do i = 1, n
+ b(i) = i
+ end do
+ !$acc end parallel loop
+
+ !$acc update host (b)
+
+ !$acc parallel loop pcopy(b) copyout(a) copyin(c)
+ do i = 1, n
+ a(i) = b(i) + c
+ end do
+ !$acc end parallel loop
+
+ !$acc exit data delete(b)
+
+ call validate
+
+ a(:) = b(:)
+ c = 0.0
+ call validate
+
+ ! Test copy
+
+ c = 1.0
+ !$acc parallel loop copy(/BLOCK/)
+ do i = 1, n
+ a(i) = b(i) + c
+ end do
+ !$acc end parallel loop
+
+ call validate
+
+ ! Test pcopyin, pcopyout FIXME
+
+ c = 2.0
+ !$acc data copyin(b, c) copyout(a)
+
+ !$acc parallel loop pcopyin(b, c) pcopyout(a)
+ do i = 1, n
+ a(i) = b(i) + c
+ end do
+ !$acc end parallel loop
+
+ !$acc end data
+
+ call validate
+
+ ! Test reduction, private
+
+ j = 0
+
+ !$acc parallel private(i) copy(j)
+ !$acc loop reduction(+:j)
+ do i = 1, n
+ j = j + 1
+ end do
+ !$acc end parallel
+
+ if (j .ne. n) call abort
+
+ ! Test firstprivate, copy
+
+ a(:) = 0
+ c = j
+
+ !$acc parallel loop firstprivate(c) copyout(a)
+ do i = 1, n
+ a(i) = i + c
+ end do
+ !$acc end parallel loop
+
+ call validate
+end program main
Index: libgomp/testsuite/libgomp.oacc-fortran/common-block-3.f90
===================================================================
--- libgomp/testsuite/libgomp.oacc-fortran/common-block-3.f90 (nonexistent)
+++ libgomp/testsuite/libgomp.oacc-fortran/common-block-3.f90 (working copy)
@@ -0,0 +1,137 @@
+! Test data located inside common blocks. This test does not execrise
+! ACC DECLARE. Most of the data clauses are implicit.
+
+module consts
+ integer, parameter :: n = 100
+end module consts
+
+subroutine validate
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 x(n), y(n), z
+ common /BLOCK/ x, y, z, j
+
+ do i = 1, n
+ if (abs(x(i) - i - z) .ge. 0.0001) call abort
+ end do
+end subroutine validate
+
+subroutine incr_parallel
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 x(n), y(n), z
+ common /BLOCK/ x, y, z, j
+
+ !$acc parallel loop
+ do i = 1, n
+ x(i) = x(i) + z
+ end do
+ !$acc end parallel loop
+end subroutine incr_parallel
+
+subroutine incr_kernels
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 x(n), y(n), z
+ common /BLOCK/ x, y, z, j
+
+ !$acc kernels
+ do i = 1, n
+ x(i) = x(i) + z
+ end do
+ !$acc end kernels
+end subroutine incr_kernels
+
+program main
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 a(n), b(n), c
+ common /BLOCK/ a, b, c, j
+
+ !$acc data copyout(a, c)
+
+ c = 1.0
+
+ !$acc update device(c)
+
+ !$acc parallel loop
+ do i = 1, n
+ a(i) = i
+ end do
+ !$acc end parallel loop
+
+ call incr_parallel
+ call incr_parallel
+ call incr_parallel
+ !$acc end data
+
+ c = 3.0
+ call validate
+
+ ! Test pcopy without copyout
+
+ c = 2.0
+ call incr_kernels
+ c = 5.0
+ call validate
+
+ !$acc kernels
+ do i = 1, n
+ b(i) = i
+ end do
+ !$acc end kernels
+
+ !$acc parallel loop
+ do i = 1, n
+ a(i) = b(i) + c
+ end do
+ !$acc end parallel loop
+
+ call validate
+
+ a(:) = b(:)
+ c = 0.0
+ call validate
+
+ ! Test copy
+
+ c = 1.0
+ !$acc parallel loop
+ do i = 1, n
+ a(i) = b(i) + c
+ end do
+ !$acc end parallel loop
+
+ call validate
+
+ c = 2.0
+ !$acc data copyin(b, c) copyout(a)
+
+ !$acc kernels
+ do i = 1, n
+ a(i) = b(i) + c
+ end do
+ !$acc end kernels
+
+ !$acc end data
+
+ call validate
+
+ j = 0
+
+ !$acc parallel loop reduction(+:j)
+ do i = 1, n
+ j = j + 1
+ end do
+ !$acc end parallel loop
+
+ if (j .ne. n) call abort
+end program main
Index: libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90
===================================================================
--- libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90 (nonexistent)
+++ libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90 (working copy)
@@ -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] 4+ messages in thread
end of thread, other threads:[~2019-01-02 14:38 UTC | newest]
Thread overview: 4+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2018-06-29 19:34 [patch] Add OpenACC Fortran support for deviceptr and variable in common blocks Cesar Philippidis
2018-12-04 13:30 ` Jakub Jelinek
2018-12-04 13:50 ` Chung-Lin Tang
2019-01-02 14:38 ` Chung-Lin Tang
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).