diff mbox series

Don't create location wrapper nodes within OpenACC clauses (was: [WIP] Fold 'NON_LVALUE_EXPR' some more (was: C++ 'NON_LVALUE_EXPR' in OMP array section handling))

Message ID 87ft4wbgt0.fsf@euler.schwinge.homeip.net
State New
Headers show
Series Don't create location wrapper nodes within OpenACC clauses (was: [WIP] Fold 'NON_LVALUE_EXPR' some more (was: C++ 'NON_LVALUE_EXPR' in OMP array section handling)) | expand

Commit Message

Thomas Schwinge Nov. 26, 2020, 9:36 a.m. UTC 
Hi!

On 2020-05-25T12:56:15+0200, I wrote:
> Anyone have any input here, especially whether something like the WIP
> patch attached to generally "Fold 'NON_LVALUE_EXPR' some more" is
> preferable over local 'STRIP_NOPS'?

Neither of these, actually...

> On 2020-03-26T20:53:19+0100, I wrote:
>> On 2020-03-26T09:09:01-0600, Sandra Loosemore <sandra@codesourcery.com> wrote:
>>> On 3/26/20 8:27 AM, Thomas Schwinge wrote:
>>>> Note that as this code is shared between OpenACC/OpenMP, this might
>>>> affect OpenMP, too, as far as I can tell.  (Subject updated.)  Jakub, can
>>>> you please have a look, too?

This is actually what made me think that something's "fishy" here: this
code path supposedly works fine for OpenMP, and now suddenly needs
adjustments for the very same OpenACC usage.

>>>> On 2020-03-25T23:02:38-0600, Sandra Loosemore <sandra@codesourcery.com> wrote:
>>>>> The attached patch fixes a bug I found in the C++ front end's handling
>>>>> of OpenACC data clauses.  The problem here is that the C++ front end
>>>>> wraps the array bounds expressions in NON_LVALUE_EXPR tree nodes, and
>>>>> the code here (which appears to have been copied from similar code in
>>>>> the C front end) was failing to strip those before checking to see if
>>>>> they were INTEGER_CST nodes, etc.
>>>>
>>>> So, I had a quick look.  I'm confirming the 'NON_LVALUE_EXPR' (C++ only,
>>>> not seen for C) difference, and that 'STRIP_NOPS' gets rid of these.
>>>> However, I also in some code paths see, for example, 'integer_nonzerop'
>>>> calls, which internally do 'STRIP_ANY_LOCATION_WRAPPER'.
>>>>
>>>> I don't know if 'STRIP_NOPS' is the correct "hammer" to use here, I don't
>>>> know what the usual convention is: explicitly remove (via 'STRIP_NOPS' as
>>>> you suggested, or something else), or have the enquiry functions do it
>>>> ('STRIP_ANY_LOCATION_WRAPPER' as 'integer_nonzerop' is doing, for
>>>> example).
>>>>
>>>> Empirical data doesn't mean too much here, of course, I'm not seeing a
>>>> lot of explicit 'STRIP_*' calls in 'gcc/cp/semantics.c'.  ;-)
>>>
>>> I saw that STRIP_NOPS seem to be the preferred way to do things in e.g.
>>> fold-const.c.  I don't know if there's a reason to use some less general
>>> form of STRIP_* here?
>
>>>>> Sadly, I have no test case for this because it was only triggering an
>>>>> error in conjunction with some other OpenACC patches that are not yet on
>>>>> trunk
>
>>> In the current code [we have]
>>> checks like
>>>
>>> TREE_CODE (low_bound) == INTEGER_CST
>>>
>>> etc.  So when they're wrapped in NON_LVALUE_EXPRs, it's basically
>>> failing to detect constants in array dimension specifiers and falling
>>> through to other code.
>>
>> Eh, indeed...  ;-\ (So we should be able to deduce some misbehavior from
>> that, to construct a test case.  I'll have a look.)
>
> (Have not yet been able to look into constructing any test cases.)

Now I have.

>> So.  I'm not objecting to handling 'NON_LVALUE_EXPR's locally here via
>> any kind of 'STRIP_*', but it somehow doesn't seem the general solution.
>> Another option seems to be to teach 'fold_simple' to handle
>> 'NON_LVALUE_EXPR's, so that the existing code:
>>
>>     /* We need to reduce to real constant-values for checks below.  */
>>     if (length)
>>       length = fold_simple (length);
>>     if (low_bound)
>>       low_bound = fold_simple (low_bound);
>>     if (low_bound
>>         && TREE_CODE (low_bound) == INTEGER_CST
>>         && [...])
>>       low_bound = fold_convert (sizetype, low_bound);
>>     if (length
>>         && TREE_CODE (length) == INTEGER_CST
>>         && [...])
>>       length = fold_convert (sizetype, length);
>>
>> ... would then just work.  But: I don't know if 'fold_simple' (and
>> others?) should handle 'NON_LVALUE_EXPR's, or if there's a reason why it
>> doesn't.  (Have not yet tried to figure that out.)  What I can tell is
>> that the attached patch does solve the issue in the C++ OMP array section
>> handling, and survives a powerpc64le-unknown-linux-gnu
>> '--enable-checking=yes' (will now re-run with 'fold' checking) bootstrap,
>> with no testsuite regressions.
>>
>> Hmm.

So, I understand, 'NON_LVALUE_EXPR's are primarily a C/C++ (only?) front
end construct to wrap nodes that must not be used as a lvalue (per the
programming language standards); rejecting code like 'int x; &(x + 0);',
I suppose.

Then, I'd assume, that after the front ends, they indeed can all be
stripped (assuming that they loose their special semantics after the
front ends, and the middle end, back ends don't (have to) care).  Thus,
it may indeed conceptually make sense to make them front end specific
nodes (and then later assert that such nodes aren't present anymore);
strip them in some "fold" functions as I had experimented/suggested:

> Subject: [PATCH] Fold 'NON_LVALUE_EXPR' some more

> --- a/gcc/cp/constexpr.c
> +++ b/gcc/cp/constexpr.c
> @@ -6650,6 +6650,7 @@ fold_simple_1 (tree t)
>      case BIT_NOT_EXPR:
>      case TRUTH_NOT_EXPR:
>      case NOP_EXPR:
> +    case NON_LVALUE_EXPR:
>      case VIEW_CONVERT_EXPR:
>      case CONVERT_EXPR:
>      case FLOAT_EXPR:
> --- a/gcc/fold-const.c
> +++ b/gcc/fold-const.c
> @@ -1739,6 +1739,7 @@ const_unop (enum tree_code code, tree type, tree arg0)
>    switch (code)
>      {
>      CASE_CONVERT:
> +    case NON_LVALUE_EXPR:
>      case FLOAT_EXPR:
>      case FIX_TRUNC_EXPR:
>      case FIXED_CONVERT_EXPR:

But: it's not clear to me which of the several "fold" functions work at
the front end level (and thus likely must preserve 'NON_LVALUE_EXPR's --
thus is it worrying that the 'gcc/cp/constexpr.c' change cited above
doesn't introduce any testsuite regressions?) and which work at the
middle end level (where stripping 'NON_LVALUE_EXPR's supposedly would be
fine, handling them like 'CASE_CONVERT': 'NOP_EXPR', 'CONVERT_EXPR').


But now, one step back: the other situation where 'NON_LVALUE_EXPR's
appear is as location wrappers (see 'gcc/tree.c:maybe_wrap_with_location'
etc.).  And this is what's the problem here.  I've just pushed "Don't
create location wrapper nodes within OpenACC clauses" to master branch in
commit c0c7270cc4efd896fe99f8ad5409dbef089a407f, and backported to
releases/gcc-10 branch in commit
e8e0357d129187b24085ce52172c87dbf6c2ecae, to releases/gcc-9 branch in
commit 25b61f935a8eca56c68c8587fc8915797250bb30, to releases/gcc-8 branch
in commit 23ec71d91e3044108a557dace573d3e60ff1c07e (testsuite changes
only), see attached.  Quoting myself from the commit log:

| This fixes a GCC 11, 10, 9 regression introduced by commit
| dfd7fdca2ac17d8b823a16700525824ca312ade0 (Subversion r267272) "C++: more
| location wrapper nodes (PR c++/43064, PR c++/43486)".  But: this isn't
| intending to blame David, because back then, the problem hasn't been visible in
| the testsuite (or else I'm sure would've been addressed right away) because of
| our all dear friend: missing testsuite coverage.  Thus, for GCC 8, I'm likewise
| enhancing the testsuite, without the C++ front end code changes.
|
| I actually had presumed that there may be an issue for OpenACC:
| <http://mid.mail-archive.com/874lb9qr2u.fsf@euler.schwinge.homeip.net>, so here
| we are, two years (and many "wasted" hours...) later...


Grüße
 Thomas


-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

Comments

Jakub Jelinek Nov. 26, 2020, 10:02 a.m. UTC | #1 
On Thu, Nov 26, 2020 at 10:36:43AM +0100, Thomas Schwinge wrote:
> So, I understand, 'NON_LVALUE_EXPR's are primarily a C/C++ (only?) front
> end construct to wrap nodes that must not be used as a lvalue (per the
> programming language standards); rejecting code like 'int x; &(x + 0);',
> I suppose.

The NON_LVALUE_EXPRs serve two purposes, one is indeed to indicate something
is not an lvalue, even when its operand is, and another is holding location
info (I think currently only in the C++ FE), so that diagnostics can report
correct locations even for e.g. constants or uses of variables.
So, it is undesirable to avoid adding those location wrappers, but instead
when we want to look at the value of the expression we should be using
appropriate APIs (e.g. for_fold_warn, or maybe_constant_value).
And generally, cp_fold_function should be also stripping them away from
everything.

	Jakub
diff mbox series

Patch

From 23ec71d91e3044108a557dace573d3e60ff1c07e Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed, 25 Nov 2020 20:36:55 +0100
Subject: [PATCH] Don't create location wrapper nodes within OpenACC clauses
 (testsuite changes only)

This fixes a GCC 11, 10, 9 regression introduced by commit
dfd7fdca2ac17d8b823a16700525824ca312ade0 (Subversion r267272) "C++: more
location wrapper nodes (PR c++/43064, PR c++/43486)".  But: this isn't
intending to blame David, because back then, the problem hasn't been visible in
the testsuite (or else I'm sure would've been addressed right away) because of
our all dear friend: missing testsuite coverage.  Thus, for GCC 8, I'm likewise
enhancing the testsuite, without the C++ front end code changes.

I actually had presumed that there may be an issue for OpenACC:
<http://mid.mail-archive.com/874lb9qr2u.fsf@euler.schwinge.homeip.net>, so here
we are, two years (and many "wasted" hours...) later...

	gcc/testsuite/
	* c-c++-common/goacc/cache-3-1.c: New.
	* c-c++-common/goacc/cache-3-2.c: Likewise.
	* c-c++-common/goacc/data-clause-1.c: Likewise.
	* c-c++-common/goacc/data-clause-2.c: Likewise.
	* c-c++-common/gomp/map-1.c: Adjust.
	* c-c++-common/gomp/map-2.c: Likewise.
	* g++.dg/goacc/cache-3-1.C: New.
	* g++.dg/goacc/cache-3-2.C: Likewise.
	* g++.dg/goacc/data-clause-1.C: Likewise.
	* g++.dg/goacc/data-clause-2.C: Likewise.
	* g++.dg/gomp/map-1.C: Adjust.
	* g++.dg/gomp/map-2.C: Likewise.

Reported-by: Sandra Loosemore <sandra@codesourcery.com>
(cherry picked from commit c0c7270cc4efd896fe99f8ad5409dbef089a407f (testsuite changes only))
---
 gcc/testsuite/c-c++-common/goacc/cache-3-1.c  | 116 +++++++++++++++++
 gcc/testsuite/c-c++-common/goacc/cache-3-2.c  |  50 +++++++
 .../c-c++-common/goacc/data-clause-1.c        | 115 ++++++++++++++++
 .../c-c++-common/goacc/data-clause-2.c        |  49 +++++++
 gcc/testsuite/c-c++-common/gomp/map-1.c       |   4 +-
 gcc/testsuite/c-c++-common/gomp/map-2.c       |   4 +-
 gcc/testsuite/g++.dg/goacc/cache-3-1.C        | 123 ++++++++++++++++++
 gcc/testsuite/g++.dg/goacc/cache-3-2.C        |  57 ++++++++
 gcc/testsuite/g++.dg/goacc/data-clause-1.C    | 122 +++++++++++++++++
 gcc/testsuite/g++.dg/goacc/data-clause-2.C    |  56 ++++++++
 gcc/testsuite/g++.dg/gomp/map-1.C             |   6 +-
 gcc/testsuite/g++.dg/gomp/map-2.C             |   4 +-
 12 files changed, 701 insertions(+), 5 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/cache-3-1.c
 create mode 100644 gcc/testsuite/c-c++-common/goacc/cache-3-2.c
 create mode 100644 gcc/testsuite/c-c++-common/goacc/data-clause-1.c
 create mode 100644 gcc/testsuite/c-c++-common/goacc/data-clause-2.c
 create mode 100644 gcc/testsuite/g++.dg/goacc/cache-3-1.C
 create mode 100644 gcc/testsuite/g++.dg/goacc/cache-3-2.C
 create mode 100644 gcc/testsuite/g++.dg/goacc/data-clause-1.C
 create mode 100644 gcc/testsuite/g++.dg/goacc/data-clause-2.C

diff --git a/gcc/testsuite/c-c++-common/goacc/cache-3-1.c b/gcc/testsuite/c-c++-common/goacc/cache-3-1.c
new file mode 100644
index 00000000000..155de689dfa
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/cache-3-1.c
@@ -0,0 +1,116 @@ 
+/* Test 'cache' directive diagnostics.  */
+
+/* See also corresponding C++ variant: '../../g++.dg/goacc/cache-3-1.C'.  */
+
+/* See also corresponding C/C++ data clause variant: 'data-clause-1.c'.  */
+
+/* { dg-additional-options "-fopenmp" } for '#pragma omp threadprivate'.  */
+
+/* The current implementation doesn't restrict where a 'cache' directive may
+   appear, so we don't make any special arrangements.  */
+
+extern int a[][10], a2[][10];
+int b[10], c[10][2], d[10], e[10], f[10];
+int b2[10], c2[10][2], d2[10], e2[10], f2[10];
+int k[10], l[10], m[10], n[10], o;
+int *p;
+int **q;
+int r[4][4][4][4][4];
+extern struct s s1;
+extern struct s s2[1]; /* { dg-error "array type has incomplete element type" "" { target c } } */
+int t[10];
+#pragma omp threadprivate (t)
+#pragma acc routine
+void bar (int *);
+
+void
+foo (int g[3][10], int h[4][8], int i[2][10], int j[][9],
+     int g2[3][10], int h2[4][8], int i2[2][10], int j2[][9])
+{
+  #pragma acc cache(bar[2:5]) /* { dg-error "is not a variable" } */
+    ;
+  #pragma acc cache(t[2:5]) /* { dg-error "is threadprivate variable" } */
+    ;
+  #pragma acc cache(k[0.5:]) /* { dg-error "low bound \[^\n\r]* of array section does not have integral type" } */
+    ;
+  #pragma acc cache(l[:7.5f]) /* { dg-error "length \[^\n\r]* of array section does not have integral type" } */
+    ;
+  #pragma acc cache(m[p:]) /* { dg-error "low bound \[^\n\r]* of array section does not have integral type" } */
+    ;
+  #pragma acc cache(n[:p]) /* { dg-error "length \[^\n\r]* of array section does not have integral type" } */
+    ;
+  #pragma acc cache(o[2:5]) /* { dg-error "does not have pointer or array type" } */
+    ;
+  #pragma acc cache(s1) /* { dg-error "expected '\\\['" } */
+    ;
+  #pragma acc cache(s2) /* { dg-error "expected '\\\['" } */
+    ;
+  #pragma acc cache(a[:][:]) /* { dg-error "array type length expression must be specified" } */
+    bar (&a[0][0]); /* { dg-bogus "referenced in target region does not have a mappable type" } */
+  #pragma acc cache(b[-1:]) /* { dg-error "negative low bound in array section" } */
+    bar (b);
+  #pragma acc cache(c[:-3][:]) /* { dg-error "negative length in array section" } */
+    bar (&c[0][0]);
+  #pragma acc cache(d[11:]) /* { dg-error "low bound \[^\n\r]* above array section size" } */
+    bar (d);
+  #pragma acc cache(e[:11]) /* { dg-error "length \[^\n\r]* above array section size" } */
+    bar (e);
+  #pragma acc cache(f[1:10]) /* { dg-error "high bound \[^\n\r]* above array section size" } */
+    bar (f);
+  #pragma acc cache(g[:][0:10]) /* { dg-error "for pointer type length expression must be specified" } */
+    bar (&g[0][0]);
+  #pragma acc cache(h[2:1][-1:]) /* { dg-error "negative low bound in array section" } */
+    bar (&h[0][0]);
+  #pragma acc cache(h[:1][:-3]) /* { dg-error "negative length in array section" } */
+    bar (&h[0][0]);
+  #pragma acc cache(i[:1][11:]) /* { dg-error "low bound \[^\n\r]* above array section size" } */
+    bar (&i[0][0]);
+  #pragma acc cache(j[3:1][:10]) /* { dg-error "length \[^\n\r]* above array section size" } */
+    bar (&j[0][0]);
+  #pragma acc cache(j[30:1][5:5]) /* { dg-error "high bound \[^\n\r]* above array section size" } */
+    bar (&j[0][0]);
+  #pragma acc cache(a2[:1][2:4])
+    bar (&a2[0][0]);
+  #pragma acc cache(a2[3:5][:])
+    bar (&a2[0][0]);
+  #pragma acc cache(a2[3:5][:10])
+    bar (&a2[0][0]);
+  #pragma acc cache(b2[0:])
+    bar (b2);
+  #pragma acc cache(c2[:3][:])
+    bar (&c2[0][0]);
+  #pragma acc cache(d2[9:])
+    bar (d2);
+  #pragma acc cache(e2[:10])
+    bar (e2);
+  #pragma acc cache(f2[1:9])
+    bar (f2);
+  #pragma acc cache(g2[:1][2:4])
+    bar (&g2[0][0]);
+  #pragma acc cache(h2[2:2][0:])
+    bar (&h2[0][0]);
+  #pragma acc cache(h2[:1][:3])
+    bar (&h2[0][0]);
+  #pragma acc cache(i2[:1][9:])
+    bar (&i2[0][0]);
+  #pragma acc cache(j2[3:4][:9])
+    bar (&j2[0][0]);
+  #pragma acc cache(j2[30:1][5:4])
+    bar (&j2[0][0]);
+  #pragma acc cache(q[1:2])
+    ;
+  #pragma acc cache(q[3:5][:10]) /* { dg-error "array section is not contiguous" } */
+    ;
+  #pragma acc cache(r[3:][2:1][1:2])
+    ;
+  #pragma acc cache(r[3:][2:1][1:2][:][0:4])
+    ;
+  #pragma acc cache(r[3:][2:1][1:2][1:][0:4]) /* { dg-error "array section is not contiguous" } */
+    ;
+  #pragma acc cache(r[3:][2:1][1:2][:3][0:4]) /* { dg-error "array section is not contiguous" } */
+    ;
+  #pragma acc cache(r[3:][2:1][1:2][:][1:]) /* { dg-error "array section is not contiguous" } */
+    ;
+  #pragma acc cache(r[3:][2:1][1:2][:][:3]) /* { dg-error "array section is not contiguous" } */
+    ;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/cache-3-2.c b/gcc/testsuite/c-c++-common/goacc/cache-3-2.c
new file mode 100644
index 00000000000..ea5222e7d0c
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/cache-3-2.c
@@ -0,0 +1,50 @@ 
+/* Test 'cache' directive diagnostics.  */
+
+/* See also corresponding C++ variant: '../../g++.dg/goacc/cache-3-2.C'.  */
+
+/* See also corresponding C/C++ data clause variant: 'data-clause-2.c'.  */
+
+/* The current implementation doesn't restrict where a 'cache' directive may
+   appear, so we don't make any special arrangements.  */
+
+void
+foo (int *p, int (*q)[10], int r[10], int s[10][10])
+{
+  int a[10], b[10][10];
+  #pragma acc cache (p[-1:2])
+  ;
+  #pragma acc cache (q[-1:2][0:10])
+  ;
+  #pragma acc cache (q[-1:2][-2:10]) /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc cache (r[-1:2])
+  ;
+  #pragma acc cache (s[-1:2][:])
+  ;
+  #pragma acc cache (s[-1:2][-2:10]) /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc cache (a[-1:2])	 /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc cache (b[-1:2][0:])	 /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc cache (b[1:2][-2:10]) /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc cache (p[2:-3])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc cache (q[2:-3][:])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc cache (q[2:3][0:-1])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc cache (r[2:-5])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc cache (s[2:-5][:])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc cache (s[2:5][0:-4])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc cache (a[2:-5])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc cache (b[2:-5][0:10]) /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc cache (b[2:5][0:-4]) /* { dg-error "negative length in array section in" } */
+  ;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/data-clause-1.c b/gcc/testsuite/c-c++-common/goacc/data-clause-1.c
new file mode 100644
index 00000000000..20f822cbb99
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/data-clause-1.c
@@ -0,0 +1,115 @@ 
+/* Test data clause diagnostics.  */
+
+/* See also corresponding OpenACC C++ variant: '../../g++.dg/goacc/data-clause-1.C'.  */
+
+/* See also corresponding OpenACC 'cache' directive variant: 'cache-3-1.c'.  */
+
+/* See also corresponding OpenMP variant: '../gomp/map-1.c'.  */
+
+/* { dg-additional-options "-fopenmp" } for '#pragma omp threadprivate'.  */
+
+extern int a[][10], a2[][10];
+int b[10], c[10][2], d[10], e[10], f[10];
+int b2[10], c2[10][2], d2[10], e2[10], f2[10];
+int k[10], l[10], m[10], n[10], o;
+int *p;
+int **q;
+int r[4][4][4][4][4];
+extern struct s s1;
+extern struct s s2[1]; /* { dg-error "array type has incomplete element type" "" { target c } } */
+int t[10];
+#pragma omp threadprivate (t)
+#pragma acc routine
+void bar (int *);
+
+void
+foo (int g[3][10], int h[4][8], int i[2][10], int j[][9],
+     int g2[3][10], int h2[4][8], int i2[2][10], int j2[][9])
+{
+  #pragma acc parallel copyin(bar[2:5]) /* { dg-error "is not a variable" } */
+    ;
+  #pragma acc parallel copyout(t[2:5]) /* { dg-error "is threadprivate variable" } */
+    ;
+  #pragma acc parallel copy(k[0.5:]) /* { dg-error "low bound \[^\n\r]* of array section does not have integral type" } */
+    ;
+  #pragma acc parallel copyout(l[:7.5f]) /* { dg-error "length \[^\n\r]* of array section does not have integral type" } */
+    ;
+  #pragma acc parallel copyin(m[p:]) /* { dg-error "low bound \[^\n\r]* of array section does not have integral type" } */
+    ;
+  #pragma acc parallel copy(n[:p]) /* { dg-error "length \[^\n\r]* of array section does not have integral type" } */
+    ;
+  #pragma acc parallel copyin(o[2:5]) /* { dg-error "does not have pointer or array type" } */
+    ;
+  #pragma acc parallel create(s1) /* { dg-error "'s1' does not have a mappable type in 'map' clause" } */
+    ;
+  #pragma acc parallel create(s2) /* { dg-error "'s2' does not have a mappable type in 'map' clause" } */
+    ;
+  #pragma acc parallel copyin(a[:][:]) /* { dg-error "array type length expression must be specified" } */
+    bar (&a[0][0]); /* { dg-error "referenced in target region does not have a mappable type" } */
+  #pragma acc parallel copy(b[-1:]) /* { dg-error "negative low bound in array section" } */
+    bar (b);
+  #pragma acc parallel copy(c[:-3][:]) /* { dg-error "negative length in array section" } */
+    bar (&c[0][0]);
+  #pragma acc parallel copyout(d[11:]) /* { dg-error "low bound \[^\n\r]* above array section size" } */
+    bar (d);
+  #pragma acc parallel copyin(e[:11]) /* { dg-error "length \[^\n\r]* above array section size" } */
+    bar (e);
+  #pragma acc parallel copyin(f[1:10]) /* { dg-error "high bound \[^\n\r]* above array section size" } */
+    bar (f);
+  #pragma acc parallel copyout(g[:][0:10]) /* { dg-error "for pointer type length expression must be specified" } */
+    bar (&g[0][0]);
+  #pragma acc parallel copyout(h[2:1][-1:]) /* { dg-error "negative low bound in array section" } */
+    bar (&h[0][0]);
+  #pragma acc parallel copy(h[:1][:-3]) /* { dg-error "negative length in array section" } */
+    bar (&h[0][0]);
+  #pragma acc parallel copy(i[:1][11:]) /* { dg-error "low bound \[^\n\r]* above array section size" } */
+    bar (&i[0][0]);
+  #pragma acc parallel copyout(j[3:1][:10]) /* { dg-error "length \[^\n\r]* above array section size" } */
+    bar (&j[0][0]);
+  #pragma acc parallel copyin(j[30:1][5:5]) /* { dg-error "high bound \[^\n\r]* above array section size" } */
+    bar (&j[0][0]);
+  #pragma acc parallel copyin(a2[:1][2:4])
+    bar (&a2[0][0]);
+  #pragma acc parallel copy(a2[3:5][:])
+    bar (&a2[0][0]);
+  #pragma acc parallel copyin(a2[3:5][:10])
+    bar (&a2[0][0]);
+  #pragma acc parallel copy(b2[0:])
+    bar (b2);
+  #pragma acc parallel copy(c2[:3][:])
+    bar (&c2[0][0]);
+  #pragma acc parallel copyout(d2[9:])
+    bar (d2);
+  #pragma acc parallel copyin(e2[:10])
+    bar (e2);
+  #pragma acc parallel copyin(f2[1:9])
+    bar (f2);
+  #pragma acc parallel copy(g2[:1][2:4])
+    bar (&g2[0][0]);
+  #pragma acc parallel copyout(h2[2:2][0:])
+    bar (&h2[0][0]);
+  #pragma acc parallel copy(h2[:1][:3])
+    bar (&h2[0][0]);
+  #pragma acc parallel copyin(i2[:1][9:])
+    bar (&i2[0][0]);
+  #pragma acc parallel copyout(j2[3:4][:9])
+    bar (&j2[0][0]);
+  #pragma acc parallel copyin(j2[30:1][5:4])
+    bar (&j2[0][0]);
+  #pragma acc parallel copy(q[1:2])
+    ;
+  #pragma acc parallel copy(q[3:5][:10]) /* { dg-error "array section is not contiguous" } */
+    ;
+  #pragma acc parallel copy(r[3:][2:1][1:2])
+    ;
+  #pragma acc parallel copy(r[3:][2:1][1:2][:][0:4])
+    ;
+  #pragma acc parallel copy(r[3:][2:1][1:2][1:][0:4]) /* { dg-error "array section is not contiguous" } */
+    ;
+  #pragma acc parallel copy(r[3:][2:1][1:2][:3][0:4]) /* { dg-error "array section is not contiguous" } */
+    ;
+  #pragma acc parallel copy(r[3:][2:1][1:2][:][1:]) /* { dg-error "array section is not contiguous" } */
+    ;
+  #pragma acc parallel copy(r[3:][2:1][1:2][:][:3]) /* { dg-error "array section is not contiguous" } */
+    ;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/data-clause-2.c b/gcc/testsuite/c-c++-common/goacc/data-clause-2.c
new file mode 100644
index 00000000000..d4603b016dd
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/data-clause-2.c
@@ -0,0 +1,49 @@ 
+/* Test data clause diagnostics.  */
+
+/* See also corresponding OpenACC C++ variant: '../../g++.dg/goacc/data-clause-2.C'.  */
+
+/* See also corresponding OpenACC 'cache' directive variant: 'cache-3-2.c'.  */
+
+/* See also corresponding OpenMP variant: '../gomp/map-2.c'.  */
+
+void
+foo (int *p, int (*q)[10], int r[10], int s[10][10])
+{
+  int a[10], b[10][10];
+  #pragma acc parallel copy (p[-1:2])
+  ;
+  #pragma acc parallel copy (q[-1:2][0:10])
+  ;
+  #pragma acc parallel copy (q[-1:2][-2:10]) /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc parallel copy (r[-1:2])
+  ;
+  #pragma acc parallel copy (s[-1:2][:])
+  ;
+  #pragma acc parallel copy (s[-1:2][-2:10]) /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc parallel copy (a[-1:2])	 /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc parallel copy (b[-1:2][0:])	 /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc parallel copy (b[1:2][-2:10]) /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc parallel copy (p[2:-3])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc parallel copy (q[2:-3][:])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc parallel copy (q[2:3][0:-1])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc parallel copy (r[2:-5])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc parallel copy (s[2:-5][:])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc parallel copy (s[2:5][0:-4])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc parallel copy (a[2:-5])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc parallel copy (b[2:-5][0:10]) /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc parallel copy (b[2:5][0:-4]) /* { dg-error "negative length in array section in" } */
+  ;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/map-1.c b/gcc/testsuite/c-c++-common/gomp/map-1.c
index cd026f1243f..a8549bc49d9 100644
--- a/gcc/testsuite/c-c++-common/gomp/map-1.c
+++ b/gcc/testsuite/c-c++-common/gomp/map-1.c
@@ -1,6 +1,8 @@ 
 /* Test 'map' clause diagnostics.  */
 
-/* See also corresponding C++ variant: '../../g++.dg/gomp/map-1.C'.  */
+/* See also corresponding OpenMP C++ variant: '../../g++.dg/gomp/map-1.C'.  */
+
+/* See also corresponding OpenACC variant: '../goacc/data-clause-1.c'.  */
 
 extern int a[][10], a2[][10];
 int b[10], c[10][2], d[10], e[10], f[10];
diff --git a/gcc/testsuite/c-c++-common/gomp/map-2.c b/gcc/testsuite/c-c++-common/gomp/map-2.c
index cd69f6b9a57..01fb4be869d 100644
--- a/gcc/testsuite/c-c++-common/gomp/map-2.c
+++ b/gcc/testsuite/c-c++-common/gomp/map-2.c
@@ -1,6 +1,8 @@ 
 /* Test 'map' clause diagnostics.  */
 
-/* See also corresponding C++ variant: '../../g++.dg/gomp/map-2.C'.  */
+/* See also corresponding OpenMP C++ variant: '../../g++.dg/gomp/map-2.C'.  */
+
+/* See also corresponding OpenACC variant: '../goacc/data-clause-2.c'.  */
 
 void
 foo (int *p, int (*q)[10], int r[10], int s[10][10])
diff --git a/gcc/testsuite/g++.dg/goacc/cache-3-1.C b/gcc/testsuite/g++.dg/goacc/cache-3-1.C
new file mode 100644
index 00000000000..d543db60a9d
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/cache-3-1.C
@@ -0,0 +1,123 @@ 
+/* Test 'cache' directive diagnostics.  */
+
+/* See also corresponding C/C++ variant: '../../c-c++-common/goacc/cache-3-1.c'.  */
+
+/* See also corresponding C++ data clause variant: 'data-clause-1.C'.  */
+
+/* { dg-additional-options "-fopenmp" } for '#pragma omp threadprivate'.  */
+
+/* The current implementation doesn't restrict where a 'cache' directive may
+   appear, so we don't make any special arrangements.  */
+
+extern int a[][10], a2[][10];
+int b[10], c[10][2], d[10], e[10], f[10];
+int b2[10], c2[10][2], d2[10], e2[10], f2[10];
+int k[10], l[10], m[10], n[10], o;
+int *p;
+int **q;
+int r[4][4][4][4][4];
+extern struct s s1;
+extern struct s s2[1]; /* { dg-error "array type has incomplete element type" "" { target c } } */
+int t[10];
+#pragma omp threadprivate (t)
+#pragma acc routine
+void bar (int *);
+
+template <int N>
+void
+foo (int g[3][10], int h[4][8], int i[2][10], int j[][9],
+     int g2[3][10], int h2[4][8], int i2[2][10], int j2[][9])
+{
+  #pragma acc cache(bar[2:5]) /* { dg-error "is not a variable" } */
+    ;
+  #pragma acc cache(t[2:5]) /* { dg-error "is threadprivate variable" } */
+    ;
+  #pragma acc cache(k[0.5:]) /* { dg-error "low bound \[^\n\r]* of array section does not have integral type" } */
+    ;
+  #pragma acc cache(l[:7.5f]) /* { dg-error "length \[^\n\r]* of array section does not have integral type" } */
+    ;
+  #pragma acc cache(m[p:]) /* { dg-error "low bound \[^\n\r]* of array section does not have integral type" } */
+    ;
+  #pragma acc cache(n[:p]) /* { dg-error "length \[^\n\r]* of array section does not have integral type" } */
+    ;
+  #pragma acc cache(o[2:5]) /* { dg-error "does not have pointer or array type" } */
+    ;
+  #pragma acc cache(s1) /* { dg-error "expected '\\\['" } */
+    ;
+  #pragma acc cache(s2) /* { dg-error "expected '\\\['" } */
+    ;
+  #pragma acc cache(a[:][:]) /* { dg-error "array type length expression must be specified" } */
+    bar (&a[0][0]);
+  #pragma acc cache(b[-1:]) /* { dg-error "negative low bound in array section" } */
+    bar (b);
+  #pragma acc cache(c[:-3][:]) /* { dg-error "negative length in array section" } */
+    bar (&c[0][0]);
+  #pragma acc cache(d[11:]) /* { dg-error "low bound \[^\n\r]* above array section size" } */
+    bar (d);
+  #pragma acc cache(e[:11]) /* { dg-error "length \[^\n\r]* above array section size" } */
+    bar (e);
+  #pragma acc cache(f[1:10]) /* { dg-error "high bound \[^\n\r]* above array section size" } */
+    bar (f);
+  #pragma acc cache(g[:][0:10]) /* { dg-error "for pointer type length expression must be specified" } */
+    bar (&g[0][0]);
+  #pragma acc cache(h[2:1][-1:]) /* { dg-error "negative low bound in array section" } */
+    bar (&h[0][0]);
+  #pragma acc cache(h[:1][:-3]) /* { dg-error "negative length in array section" } */
+    bar (&h[0][0]);
+  #pragma acc cache(i[:1][11:]) /* { dg-error "low bound \[^\n\r]* above array section size" } */
+    bar (&i[0][0]);
+  #pragma acc cache(j[3:1][:10]) /* { dg-error "length \[^\n\r]* above array section size" } */
+    bar (&j[0][0]);
+  #pragma acc cache(j[30:1][5:5]) /* { dg-error "high bound \[^\n\r]* above array section size" } */
+    bar (&j[0][0]);
+  #pragma acc cache(a2[:1][2:4])
+    bar (&a2[0][0]);
+  #pragma acc cache(a2[3:5][:])
+    bar (&a2[0][0]);
+  #pragma acc cache(a2[3:5][:10])
+    bar (&a2[0][0]);
+  #pragma acc cache(b2[0:])
+    bar (b2);
+  #pragma acc cache(c2[:3][:])
+    bar (&c2[0][0]);
+  #pragma acc cache(d2[9:])
+    bar (d2);
+  #pragma acc cache(e2[:10])
+    bar (e2);
+  #pragma acc cache(f2[1:9])
+    bar (f2);
+  #pragma acc cache(g2[:1][2:4])
+    bar (&g2[0][0]);
+  #pragma acc cache(h2[2:2][0:])
+    bar (&h2[0][0]);
+  #pragma acc cache(h2[:1][:3])
+    bar (&h2[0][0]);
+  #pragma acc cache(i2[:1][9:])
+    bar (&i2[0][0]);
+  #pragma acc cache(j2[3:4][:9])
+    bar (&j2[0][0]);
+  #pragma acc cache(j2[30:1][5:4])
+    bar (&j2[0][0]);
+  #pragma acc cache(q[1:2])
+    ;
+  #pragma acc cache(q[3:5][:10]) /* { dg-error "array section is not contiguous" } */
+    ;
+  #pragma acc cache(r[3:][2:1][1:2])
+    ;
+  #pragma acc cache(r[3:][2:1][1:2][:][0:4])
+    ;
+  #pragma acc cache(r[3:][2:1][1:2][1:][0:4]) /* { dg-error "array section is not contiguous" } */
+    ;
+  #pragma acc cache(r[3:][2:1][1:2][:3][0:4]) /* { dg-error "array section is not contiguous" } */
+    ;
+  #pragma acc cache(r[3:][2:1][1:2][:][1:]) /* { dg-error "array section is not contiguous" } */
+    ;
+  #pragma acc cache(r[3:][2:1][1:2][:][:3]) /* { dg-error "array section is not contiguous" } */
+    ;
+}
+
+static void
+instantiate ()
+{
+  &foo<0>;
+}
diff --git a/gcc/testsuite/g++.dg/goacc/cache-3-2.C b/gcc/testsuite/g++.dg/goacc/cache-3-2.C
new file mode 100644
index 00000000000..5561e176a56
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/cache-3-2.C
@@ -0,0 +1,57 @@ 
+/* Test 'cache' directive diagnostics.  */
+
+/* See also corresponding C/C++ variant: '../../c-c++-common/goacc/cache-3-2.c'.  */
+
+/* See also corresponding C++ data clause variant: 'data-clause-2.C'.  */
+
+/* The current implementation doesn't restrict where a 'cache' directive may
+   appear, so we don't make any special arrangements.  */
+
+template <int N>
+void
+foo (int *p, int (*q)[10], int r[10], int s[10][10])
+{
+  int a[10], b[10][10];
+  #pragma acc cache (p[-1:2])
+  ;
+  #pragma acc cache (q[-1:2][0:10])
+  ;
+  #pragma acc cache (q[-1:2][-2:10]) /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc cache (r[-1:2])
+  ;
+  #pragma acc cache (s[-1:2][:])
+  ;
+  #pragma acc cache (s[-1:2][-2:10]) /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc cache (a[-1:2])	 /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc cache (b[-1:2][0:])	 /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc cache (b[1:2][-2:10]) /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc cache (p[2:-3])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc cache (q[2:-3][:])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc cache (q[2:3][0:-1])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc cache (r[2:-5])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc cache (s[2:-5][:])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc cache (s[2:5][0:-4])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc cache (a[2:-5])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc cache (b[2:-5][0:10]) /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc cache (b[2:5][0:-4]) /* { dg-error "negative length in array section in" } */
+  ;
+}
+
+static void
+instantiate ()
+{
+  &foo<0>;
+}
diff --git a/gcc/testsuite/g++.dg/goacc/data-clause-1.C b/gcc/testsuite/g++.dg/goacc/data-clause-1.C
new file mode 100644
index 00000000000..3de3834efae
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/data-clause-1.C
@@ -0,0 +1,122 @@ 
+/* Test data clause diagnostics.  */
+
+/* See also corresponding OpenACC C/C++ variant: '../../c-c++-common/goacc/data-clause-1.c'.  */
+
+/* See also corresponding OpenACC 'cache' directive variant: 'cache-3-1.C'.  */
+
+/* See also corresponding OpenMP variant: '../gomp/map-1.C'.  */
+
+/* { dg-additional-options "-fopenmp" } for '#pragma omp threadprivate'.  */
+
+extern int a[][10], a2[][10];
+int b[10], c[10][2], d[10], e[10], f[10];
+int b2[10], c2[10][2], d2[10], e2[10], f2[10];
+int k[10], l[10], m[10], n[10], o;
+int *p;
+int **q;
+int r[4][4][4][4][4];
+extern struct s s1;
+extern struct s s2[1]; /* { dg-error "array type has incomplete element type" "" { target c } } */
+int t[10];
+#pragma omp threadprivate (t)
+#pragma acc routine
+void bar (int *);
+
+template <int N>
+void
+foo (int g[3][10], int h[4][8], int i[2][10], int j[][9],
+     int g2[3][10], int h2[4][8], int i2[2][10], int j2[][9])
+{
+  #pragma acc parallel copyin(bar[2:5]) /* { dg-error "is not a variable" } */
+    ;
+  #pragma acc parallel copyout(t[2:5]) /* { dg-error "is threadprivate variable" } */
+    ;
+  #pragma acc parallel copy(k[0.5:]) /* { dg-error "low bound \[^\n\r]* of array section does not have integral type" } */
+    ;
+  #pragma acc parallel copyout(l[:7.5f]) /* { dg-error "length \[^\n\r]* of array section does not have integral type" } */
+    ;
+  #pragma acc parallel copyin(m[p:]) /* { dg-error "low bound \[^\n\r]* of array section does not have integral type" } */
+    ;
+  #pragma acc parallel copy(n[:p]) /* { dg-error "length \[^\n\r]* of array section does not have integral type" } */
+    ;
+  #pragma acc parallel copyin(o[2:5]) /* { dg-error "does not have pointer or array type" } */
+    ;
+  #pragma acc parallel create(s1) /* { dg-error "'s1' does not have a mappable type in 'map' clause" } */
+    ;
+  #pragma acc parallel create(s2) /* { dg-error "'s2' does not have a mappable type in 'map' clause" } */
+    ;
+  #pragma acc parallel copyin(a[:][:]) /* { dg-error "array type length expression must be specified" } */
+    bar (&a[0][0]); /* { dg-error "referenced in target region does not have a mappable type" "PR97996" { xfail *-*-* } } */
+  #pragma acc parallel copy(b[-1:]) /* { dg-error "negative low bound in array section" } */
+    bar (b);
+  #pragma acc parallel copy(c[:-3][:]) /* { dg-error "negative length in array section" } */
+    bar (&c[0][0]);
+  #pragma acc parallel copyout(d[11:]) /* { dg-error "low bound \[^\n\r]* above array section size" } */
+    bar (d);
+  #pragma acc parallel copyin(e[:11]) /* { dg-error "length \[^\n\r]* above array section size" } */
+    bar (e);
+  #pragma acc parallel copyin(f[1:10]) /* { dg-error "high bound \[^\n\r]* above array section size" } */
+    bar (f);
+  #pragma acc parallel copyout(g[:][0:10]) /* { dg-error "for pointer type length expression must be specified" } */
+    bar (&g[0][0]);
+  #pragma acc parallel copyout(h[2:1][-1:]) /* { dg-error "negative low bound in array section" } */
+    bar (&h[0][0]);
+  #pragma acc parallel copy(h[:1][:-3]) /* { dg-error "negative length in array section" } */
+    bar (&h[0][0]);
+  #pragma acc parallel copy(i[:1][11:]) /* { dg-error "low bound \[^\n\r]* above array section size" } */
+    bar (&i[0][0]);
+  #pragma acc parallel copyout(j[3:1][:10]) /* { dg-error "length \[^\n\r]* above array section size" } */
+    bar (&j[0][0]);
+  #pragma acc parallel copyin(j[30:1][5:5]) /* { dg-error "high bound \[^\n\r]* above array section size" } */
+    bar (&j[0][0]);
+  #pragma acc parallel copyin(a2[:1][2:4])
+    bar (&a2[0][0]);
+  #pragma acc parallel copy(a2[3:5][:])
+    bar (&a2[0][0]);
+  #pragma acc parallel copyin(a2[3:5][:10])
+    bar (&a2[0][0]);
+  #pragma acc parallel copy(b2[0:])
+    bar (b2);
+  #pragma acc parallel copy(c2[:3][:])
+    bar (&c2[0][0]);
+  #pragma acc parallel copyout(d2[9:])
+    bar (d2);
+  #pragma acc parallel copyin(e2[:10])
+    bar (e2);
+  #pragma acc parallel copyin(f2[1:9])
+    bar (f2);
+  #pragma acc parallel copy(g2[:1][2:4])
+    bar (&g2[0][0]);
+  #pragma acc parallel copyout(h2[2:2][0:])
+    bar (&h2[0][0]);
+  #pragma acc parallel copy(h2[:1][:3])
+    bar (&h2[0][0]);
+  #pragma acc parallel copyin(i2[:1][9:])
+    bar (&i2[0][0]);
+  #pragma acc parallel copyout(j2[3:4][:9])
+    bar (&j2[0][0]);
+  #pragma acc parallel copyin(j2[30:1][5:4])
+    bar (&j2[0][0]);
+  #pragma acc parallel copy(q[1:2])
+    ;
+  #pragma acc parallel copy(q[3:5][:10]) /* { dg-error "array section is not contiguous" } */
+    ;
+  #pragma acc parallel copy(r[3:][2:1][1:2])
+    ;
+  #pragma acc parallel copy(r[3:][2:1][1:2][:][0:4])
+    ;
+  #pragma acc parallel copy(r[3:][2:1][1:2][1:][0:4]) /* { dg-error "array section is not contiguous" } */
+    ;
+  #pragma acc parallel copy(r[3:][2:1][1:2][:3][0:4]) /* { dg-error "array section is not contiguous" } */
+    ;
+  #pragma acc parallel copy(r[3:][2:1][1:2][:][1:]) /* { dg-error "array section is not contiguous" } */
+    ;
+  #pragma acc parallel copy(r[3:][2:1][1:2][:][:3]) /* { dg-error "array section is not contiguous" } */
+    ;
+}
+
+static void
+instantiate ()
+{
+  &foo<0>;
+}
diff --git a/gcc/testsuite/g++.dg/goacc/data-clause-2.C b/gcc/testsuite/g++.dg/goacc/data-clause-2.C
new file mode 100644
index 00000000000..57d1823aede
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/data-clause-2.C
@@ -0,0 +1,56 @@ 
+/* Test data clause diagnostics.  */
+
+/* See also corresponding OpenACC C/C++ variant: '../../c-c++-common/goacc/data-clause-2.c'.  */
+
+/* See also corresponding OpenACC 'cache' directive variant: 'cache-3-2.C'.  */
+
+/* See also corresponding OpenMP variant: '../gomp/map-2.C'.  */
+
+template <int N>
+void
+foo (int *p, int (*q)[10], int r[10], int s[10][10])
+{
+  int a[10], b[10][10];
+  #pragma acc parallel copy (p[-1:2])
+  ;
+  #pragma acc parallel copy (q[-1:2][0:10])
+  ;
+  #pragma acc parallel copy (q[-1:2][-2:10]) /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc parallel copy (r[-1:2])
+  ;
+  #pragma acc parallel copy (s[-1:2][:])
+  ;
+  #pragma acc parallel copy (s[-1:2][-2:10]) /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc parallel copy (a[-1:2])	 /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc parallel copy (b[-1:2][0:])	 /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc parallel copy (b[1:2][-2:10]) /* { dg-error "negative low bound in array section in" } */
+  ;
+  #pragma acc parallel copy (p[2:-3])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc parallel copy (q[2:-3][:])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc parallel copy (q[2:3][0:-1])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc parallel copy (r[2:-5])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc parallel copy (s[2:-5][:])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc parallel copy (s[2:5][0:-4])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc parallel copy (a[2:-5])	 /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc parallel copy (b[2:-5][0:10]) /* { dg-error "negative length in array section in" } */
+  ;
+  #pragma acc parallel copy (b[2:5][0:-4]) /* { dg-error "negative length in array section in" } */
+  ;
+}
+
+static void
+instantiate ()
+{
+  &foo<0>;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/map-1.C b/gcc/testsuite/g++.dg/gomp/map-1.C
index 107b59ee87a..c314d55a621 100644
--- a/gcc/testsuite/g++.dg/gomp/map-1.C
+++ b/gcc/testsuite/g++.dg/gomp/map-1.C
@@ -1,6 +1,8 @@ 
 /* Test 'map' clause diagnostics.  */
 
-/* See also corresponding C/C++ variant: '../../c-c++-common/gomp/map-1.c'.  */
+/* See also corresponding OpenMP C/C++ variant: '../../c-c++-common/gomp/map-1.c'.  */
+
+/* See also corresponding OpenACC variant: '../goacc/data-clause-1.C'.  */
 
 extern int a[][10], a2[][10];
 int b[10], c[10][2], d[10], e[10], f[10];
@@ -41,7 +43,7 @@  foo (int g[3][10], int h[4][8], int i[2][10], int j[][9],
   #pragma omp target map(alloc: s2) /* { dg-error "'s2' does not have a mappable type in 'map' clause" } */
     ;
   #pragma omp target map(to: a[:][:]) /* { dg-error "array type length expression must be specified" } */
-    bar (&a[0][0]); /* { dg-error "referenced in target region does not have a mappable type" "TODO" { xfail *-*-* } } */
+    bar (&a[0][0]); /* { dg-error "referenced in target region does not have a mappable type" "PR97996" { xfail *-*-* } } */
   #pragma omp target map(tofrom: b[-1:]) /* { dg-error "negative low bound in array section" } */
     bar (b);
   #pragma omp target map(tofrom: c[:-3][:]) /* { dg-error "negative length in array section" } */
diff --git a/gcc/testsuite/g++.dg/gomp/map-2.C b/gcc/testsuite/g++.dg/gomp/map-2.C
index 10eaaa948b8..bbe26061fe3 100644
--- a/gcc/testsuite/g++.dg/gomp/map-2.C
+++ b/gcc/testsuite/g++.dg/gomp/map-2.C
@@ -1,6 +1,8 @@ 
 /* Test 'map' clause diagnostics.  */
 
-/* See also corresponding C/C++ variant: '../../c-c++-common/gomp/map-2.c'.  */
+/* See also corresponding OpenMP C/C++ variant: '../../c-c++-common/gomp/map-2.c'.  */
+
+/* See also corresponding OpenACC variant: '../goacc/data-clause-2.C'.  */
 
 template <int N>
 void
-- 
2.17.1