Hi Thomas,
On 10/25/19 10:43 AM, Thomas Schwinge wrote:
OK for trunk, with the following few small items considered.
Committed as Rev. 277451 – after a fresh bootstrap and regtesting.
Changes:
* I have now a new test case
libgomp/testsuite/libgomp.oacc-fortran/common-block-3.f90 which looks at
omplower.
* In the compile-time *{2,3} test case, there is now also a 'enter data'
and 'update host/self/device' test.
* the libgomp tests have a 'dg-do run'.
* I modified the code in gimplify.c as proposed.
Regarding the new test case: Without the gcc/gimplify.c changes, one has
(see last item before child fn):
#pragma omp target oacc_parallel map(tofrom:a [len: 400])
map(tofrom:b [len: 400]) map(tofrom:c [len: 4]) map(tofrom:block [len:
812]) [child fn …
#pragma omp target oacc_kernels map(force_tofrom:i [len: 4])
map(tofrom:y [len: 400]) map(tofrom:x [len: 400])
map(tofrom:kernel_block [len: 804]) map(force_tofrom:c [len: 4])
map(tofrom:block [len: 812]) [child fn …
With the changes of gcc/gimplify.c, one has:
#pragma omp target oacc_parallel map(tofrom:a [len: 400])
map(tofrom:b [len: 400]) map(tofrom:c [len: 4]) [child fn …
#pragma omp target oacc_kernels map(force_tofrom:i [len: 4])
map(tofrom:y [len: 400]) map(tofrom:x [len: 400]) map(force_tofrom:c
[len: 4]) [child fn …
And without gimplify.c, the added run-tests indeed fail with:
libgomp: Trying to map into device [0x407100..0x407294) object when
[0x407100..0x407290) is already mapped
Tobias
PS:
Or, would it be easy to add an OpenACC 'kernels' test case that otherwise
faild (at run time, say, with aforementioned duplicate mapping errors, or
would contain "strange"/duplicate/conflicting mapping items in the
'-fdump-tree-gimple' dump)?
See new test case and result for the current tests.
Additionally, I have applied:
Wouldn't it be clearer if that latter one were written as follows:
if (DECL_HAS_VALUE_EXPR_P (decl))
{
if (ctx->region_type & ORT_ACC)
/* For OpenACC, defer expansion of value to avoid transfering
privatized common block data instead of im-/explicitly transfered
variables which are in common blocks. */
;
else
{
tree value = get_base_address (DECL_VALUE_EXPR (decl));
if (value && DECL_P (value) && DECL_THREAD_LOCAL_P (value))
return omp_notice_threadprivate_variable (ctx, decl, value);
}
}
@@ -7353,7 +7374,9 @@ 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);
+ /* For OpenACC, as remarked above, defer expansion. */
+ shared = !(ctx->region_type & ORT_ACC);
+ ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
Also more explicit, easier to read:
if (ctx->region_type & ORT_ACC)
/* For OpenACC, as remarked above, defer expansion. */
shared = false;
else
shared = true;
@@ -7521,6 +7544,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree
decl, bool in_code)
}
shared = ((flags | n->value) & GOVD_SHARED) != 0;
+ /* For OpenACC, cf. remark above regaring common blocks. */
+ if (ctx->region_type & ORT_ACC)
+ shared = false;
ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
And again:
if (ctx->region_type & ORT_ACC)
/* For OpenACC, cf. remark above regaring common blocks. */
shared = false;
else
shared = ((flags | n->value) & GOVD_SHARED) != 0;
(In all three cases, using an easy 'if (ctx->region_type & ORT_ACC)' to
point out the special case.)
It's still some kind of voodoo to me -- but at least, you've now also
reviewed this, and it's now documented what's going on.
And changed the test case based on:
+ !$acc exit data delete(/blockA/, /blockB/, e, v)
I note there is one single 'exit data' test, but no 'enter data'.
Also, 'update' is missing, to test the 'device' and 'self'/'host' clauses.
+ !$acc data deviceptr(/blockA/, /blockB/, e, v) ! { dg-error "Syntax error in
OpenMP variable list" }
+ !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" }
+
+ !$acc data deviceptr(/blockA/, /blockB/, e, v) ! { dg-error "Syntax error in
OpenMP variable list" }
+ !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" }
+end program test
Is there a reason for the duplicated 'deviceptr' testing?
Move 'data deviceptr' up a little bit, next to the other 'data' construct
testing?
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
Similarly.
commit 96d1e6235a5b7c81df7940c1c8727f87dc1b577a
Author: burnus <burnus@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Fri Oct 25 14:28:40 2019 +0000
[Fortran] OpenACC – permit common blocks in some clauses
2019-10-25 Cesar Philippidis <ce...@codesourcery.com>
Tobias Burnus <tob...@codesourcery.com>
gcc/fortran/
* openmp.c (gfc_match_omp_map_clause): Add and pass allow_commons
argument.
(gfc_match_omp_clauses): Update calls to permit common blocks for
OpenACC's copy/copyin/copyout, create/delete, host,
pcopy/pcopy_in/pcopy_out, present_or_copy, present_or_copy_in,
present_or_copy_out, present_or_create and self.
gcc/
* gimplify.c (oacc_default_clause): Privatize fortran common blocks.
(omp_notice_variable): Defer the expansion of DECL_VALUE_EXPR for
common block decls.
gcc/testsuite/
* gfortran.dg/goacc/common-block-1.f90: New test.
* gfortran.dg/goacc/common-block-2.f90: New test.
* gfortran.dg/goacc/common-block-3.f90: New test.
libgomp/
* 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.
Reviewed-by: Thomas Schwinge <tho...@codesourcery.com>
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@277451 138bc75d-0d04-0410-961f-82ee72b054a4
diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index cef0a3f34b6..1da576b5468 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,10 @@
+2019-10-25 Cesar Philippidis <ce...@codesourcery.com>
+ Tobias Burnus <tob...@codesourcery.com>
+
+ * gimplify.c (oacc_default_clause): Privatize fortran common blocks.
+ (omp_notice_variable): Defer the expansion of DECL_VALUE_EXPR for
+ common block decls.
+
2019-10-25 Richard Biener <rguent...@suse.de>
PR tree-optimization/92222
diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog
index 4e3910298b6..d14d190b0bd 100644
--- a/gcc/fortran/ChangeLog
+++ b/gcc/fortran/ChangeLog
@@ -1,3 +1,13 @@
+2019-10-25 Cesar Philippidis <ce...@codesourcery.com>
+ Tobias Burnus <tob...@codesourcery.com>
+
+ * openmp.c (gfc_match_omp_map_clause): Add and pass allow_commons
+ argument.
+ (gfc_match_omp_clauses): Update calls to permit common blocks for
+ OpenACC's copy/copyin/copyout, create/delete, host,
+ pcopy/pcopy_in/pcopy_out, present_or_copy, present_or_copy_in,
+ present_or_copy_out, present_or_create and self.
+
2019-10-24 Martin Liska <mli...@suse.cz>
PR fortran/92174
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 5c91fcdfd31..ca342788545 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -926,10 +926,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 allow_common)
{
gfc_omp_namelist **head = NULL;
- if (gfc_match_omp_variable_list ("", list, false, NULL, &head, true)
+ if (gfc_match_omp_variable_list ("", list, allow_common, NULL, &head, true)
== MATCH_YES)
{
gfc_omp_namelist *n;
@@ -1051,7 +1052,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, true))
continue;
if (mask & OMP_CLAUSE_COPYIN)
{
@@ -1059,7 +1060,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 (",
@@ -1070,7 +1071,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 (",
@@ -1080,7 +1081,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':
@@ -1116,7 +1117,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)
@@ -1168,12 +1169,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, true))
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
@@ -1251,7 +1252,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':
@@ -1523,47 +1524,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
@@ -1781,7 +1782,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
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 05ae2f1552b..fdf6b695003 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -7227,15 +7227,28 @@ 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);
+ /* For Fortran COMMON blocks, only used variables in those blocks are
+ transfered and remapped. The block itself will have a private clause to
+ avoid transfering the data twice.
+ The hook evaluates to false by default. For a variable in Fortran's COMMON
+ or EQUIVALENCE block, returns 'true' (as we have shared=false) - as only
+ the variables in such a COMMON/EQUIVALENCE block shall be privatized not
+ the whole block. For C++ and Fortran, it can also be true under certain
+ other conditions, if DECL_HAS_VALUE_EXPR. */
+ 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;
@@ -7246,7 +7259,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_FIRSTPRIVATE;
+ else if (AGGREGATE_TYPE_P (type))
{
/* Aggregates default to 'present_or_copy', or 'present'. */
if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
@@ -7263,7 +7278,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))
{
@@ -7327,10 +7344,18 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
if (DECL_HAS_VALUE_EXPR_P (decl))
{
- tree value = get_base_address (DECL_VALUE_EXPR (decl));
+ if (ctx->region_type & ORT_ACC)
+ /* For OpenACC, defer expansion of value to avoid transfering
+ privatized common block data instead of im-/explicitly transfered
+ variables which are in common blocks. */
+ ;
+ else
+ {
+ tree value = get_base_address (DECL_VALUE_EXPR (decl));
- if (value && DECL_P (value) && DECL_THREAD_LOCAL_P (value))
- return omp_notice_threadprivate_variable (ctx, decl, value);
+ if (value && DECL_P (value) && DECL_THREAD_LOCAL_P (value))
+ return omp_notice_threadprivate_variable (ctx, decl, value);
+ }
}
if (gimplify_omp_ctxp->outer_context == NULL
@@ -7361,7 +7386,13 @@ 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);
+ if (ctx->region_type & ORT_ACC)
+ /* For OpenACC, as remarked above, defer expansion. */
+ shared = false;
+ else
+ shared = true;
+
+ ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
if (n == NULL)
{
unsigned nflags = flags;
@@ -7528,7 +7559,11 @@ 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)
+ /* For OpenACC, as remarked above, defer expansion. */
+ shared = false;
+ else
+ shared = ((flags | n->value) & GOVD_SHARED) != 0;
ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
/* If nothing changed, there's nothing left to do. */
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 0b25ce9717e..ddf575ba8c3 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,10 @@
+2019-10-25 Cesar Philippidis <ce...@codesourcery.com>
+ Tobias Burnus <tob...@codesourcery.com>
+
+ * gfortran.dg/goacc/common-block-1.f90: New test.
+ * gfortran.dg/goacc/common-block-2.f90: New test.
+ * gfortran.dg/goacc/common-block-3.f90: New test.
+
2019-10-25 David Edelsohn <dje....@gmail.com>
* gcc.target/powerpc/pr70100.c: Add -mvsx.
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..ea437526b46
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
@@ -0,0 +1,74 @@
+! Test data clauses involving common blocks and common block data.
+! Specifically, validates 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 parallel private(/blockA/, /blockB/, e, v)
+ !$acc end parallel
+
+ !$acc parallel firstprivate(/blockA/, /blockB/, e, v)
+ !$acc update device(/blockA/)
+ !$acc update self(/blockB/, v)
+ !$acc update host(/blockA/, e, /blockB/)
+ !$acc end parallel
+
+ !$acc enter data pcopyin(/blockA/, /blockB/, e, v)
+ !$acc exit data delete(/blockA/, /blockB/, e, v)
+
+
+ ! No /block/ permitted in present and deviceptr:
+
+ !$acc data present(/blockA/, /blockB/, e, v) ! { dg-error "Syntax error in OpenMP variable list" }
+ !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" }
+
+ !$acc data deviceptr(/blockA/, /blockB/, e, v) ! { dg-error "Syntax error in OpenMP variable list" }
+ !$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..1ba945019f9
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
@@ -0,0 +1,53 @@
+! 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 update device(b, /blockA/, x) ! { dg-error "Symbol .x. present on multiple clauses" }
+ !$acc update self(z, /blockB/, v) ! { dg-error "Symbol .z. present on multiple clauses" }
+ !$acc update host(/blockA/, c) ! { dg-error "Symbol .c. present on multiple clauses" }
+ !$acc end parallel
+
+ !$acc enter data copyin(/blockB/, e, v, a, c, y) ! { dg-error "Symbol .y. present on multiple clauses" }
+ !$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/common-block-3.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
new file mode 100644
index 00000000000..9032d9331f0
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
@@ -0,0 +1,39 @@
+! { dg-options "-fopenacc -fdump-tree-omplower" }
+
+module consts
+ integer, parameter :: n = 100
+end module consts
+
+program main
+ use consts
+ implicit none
+
+ integer :: i, j
+ real :: a(n) = 0, b(n) = 0, c, d
+ real :: x(n) = 0, y(n), z
+ common /BLOCK/ a, b, c, j, d
+ common /KERNELS_BLOCK/ x, y, z
+
+ c = 1.0
+ !$acc parallel loop copy(/BLOCK/)
+ do i = 1, n
+ a(i) = b(i) + c
+ end do
+ !$acc kernels
+ do i = 1, n
+ x(i) = y(i) + c
+ end do
+ !$acc end kernels
+end program main
+
+! { dg-final { scan-tree-dump-times "omp target oacc_parallel .*map\\(tofrom:a \\\[len: 400\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_parallel .*map\\(tofrom:b \\\[len: 400\\\]\\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_parallel .*map\\(tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } }
+
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:i \\\[len: 4\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(tofrom:x \\\[len: 400\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(tofrom:y \\\[len: 400\\\]\\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } }
+
+! { dg-final { scan-tree-dump-not "map\\(.*:block\\)" "omplower" } }
+! { dg-final { scan-tree-dump-not "map\\(.*:kernels_block\\)" "omplower" } }
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 62a18ad2882..351df1153fd 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,10 @@
+2019-10-25 Cesar Philippidis <ce...@codesourcery.com>
+ Tobias Burnus <tob...@codesourcery.com>
+
+ * 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.
+
2019-10-14 Jakub Jelinek <ja...@redhat.com>
PR libgomp/92081
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..000d811a059
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/common-block-1.f90
@@ -0,0 +1,107 @@
+! { dg-do run }
+!
+! Test data located inside common blocks. This test does not exercise
+! 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..4cfcded244d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/common-block-2.f90
@@ -0,0 +1,152 @@
+! { dg-do run }
+!
+! Test data located inside common blocks. This test does not exercise
+! 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..5a68b485b1e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/common-block-3.f90
@@ -0,0 +1,139 @@
+! { dg-do run }
+!
+! Test data located inside common blocks. This test does not exercise
+! 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
