Hi! On Tue, 14 Jan 2014 16:10:05 +0100, I wrote: > --- gcc/gimplify.c > +++ gcc/gimplify.c > @@ -69,7 +69,13 @@ enum gimplify_omp_var_data
> + /* Force a specific behavior (or else, a run-time error). */
> + GOVD_MAP_FORCE = 16384,
> @@ -86,7 +92,11 @@ enum omp_region_type
> + /* Default to GOVD_MAP_FORCE for implicit mappings in this region. */
> + ORT_TARGET_MAP_FORCE = 64
> };
> @@ -6135,9 +6197,14 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void
> *data)
> OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1;
> else if (code == OMP_CLAUSE_MAP)
> {
> - OMP_CLAUSE_MAP_KIND (clause) = flags & GOVD_MAP_TO_ONLY
> - ? OMP_CLAUSE_MAP_TO
> - : OMP_CLAUSE_MAP_TOFROM;
> + unsigned map_kind;
> + map_kind = (flags & GOVD_MAP_TO_ONLY
> + ? OMP_CLAUSE_MAP_TO
> + : OMP_CLAUSE_MAP_TOFROM);
> + if (flags & GOVD_MAP_FORCE)
> + map_kind |= OMP_CLAUSE_MAP_FORCE;
> + OMP_CLAUSE_MAP_KIND (clause) = (enum omp_clause_map_kind) map_kind;
> +
> if (DECL_SIZE (decl)
> && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
> {
> @@ -6389,9 +6456,10 @@ gimplify_oacc_parallel (tree *expr_p, gimple_seq
> *pre_p)
> tree expr = *expr_p;
> gimple g;
> gimple_seq body = NULL;
> + enum omp_region_type ort =
> + (enum omp_region_type) (ORT_TARGET | ORT_TARGET_MAP_FORCE);
>
> - gimplify_scan_omp_clauses (&OACC_PARALLEL_CLAUSES (expr), pre_p,
> - ORT_TARGET);
> + gimplify_scan_omp_clauses (&OACC_PARALLEL_CLAUSES (expr), pre_p, ort);
>
> push_gimplify_context ();
I don't remember what I have been thinking when implementing this -- per
the OpenACC specification's rules for implicitly determined data
attributes, it should be present_or_copy (that is, OpenMP's tofrom,
without "force" semantics), and firstprivate/copy for scalar variables
for the parallel/kernels constructs, respectively (which is still to be
implemented, for now not considering scalar variables different from
non-scalar ones). Committed to gomp-4_0-branch in r217482:
commit 7058203891bd6e1696763603673090f161e172b8
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Thu Nov 13 12:18:34 2014 +0000
Middle end: Don't use mapping kinds with "force" semantics for OpenACC.
..., which is the wrong thing to do. Also extend libgomp to actually
distinguish between "non-force"/"force" semantics.
gcc/
* gimplify.c (gimplify_omp_workshare) <OACC_DATA, OACC_KERNELS,
OACC_PARALLEL>: Don't request ORT_TARGET_MAP_FORCE.
(enum gimplify_omp_var_data, enum omp_region_type): Remove
GOVD_MAP_FORCE, and ORT_TARGET_MAP_FORCE, respectively. Update
all users.
include/
* gomp-constants.h: Define _GOMP_MAP_FLAG_SPECIAL and
_GOMP_MAP_FLAG_FORCE.
libgomp/
* target.c (gomp_map_vars_existing): Error out if "force"
semantics.
(gomp_map_vars): Actually pass kinds to gomp_map_vars_existing.
Remove FIXMEs.
* testsuite/libgomp.oacc-c-c++-common/data-already-1.c: New file.
* testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-4.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-5.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-6.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-7.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-1.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-2.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-3.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-4.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-5.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-6.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-7.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-8.f: Likewise.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@217482
138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 6 ++
gcc/gimplify.c | 65 +++-------------------
include/ChangeLog.gomp | 4 ++
include/gomp-constants.h | 3 +
libgomp/ChangeLog.gomp | 23 ++++++++
libgomp/target.c | 21 ++++---
.../libgomp.oacc-c-c++-common/data-already-1.c | 19 +++++++
.../libgomp.oacc-c-c++-common/data-already-2.c | 16 ++++++
.../libgomp.oacc-c-c++-common/data-already-3.c | 17 ++++++
.../libgomp.oacc-c-c++-common/data-already-4.c | 17 ++++++
.../libgomp.oacc-c-c++-common/data-already-5.c | 17 ++++++
.../libgomp.oacc-c-c++-common/data-already-6.c | 17 ++++++
.../libgomp.oacc-c-c++-common/data-already-7.c | 17 ++++++
.../libgomp.oacc-c-c++-common/data-already-8.c | 16 ++++++
.../libgomp.oacc-fortran/data-already-1.f | 17 ++++++
.../libgomp.oacc-fortran/data-already-2.f | 16 ++++++
.../libgomp.oacc-fortran/data-already-3.f | 15 +++++
.../libgomp.oacc-fortran/data-already-4.f | 14 +++++
.../libgomp.oacc-fortran/data-already-5.f | 14 +++++
.../libgomp.oacc-fortran/data-already-6.f | 14 +++++
.../libgomp.oacc-fortran/data-already-7.f | 14 +++++
.../libgomp.oacc-fortran/data-already-8.f | 16 ++++++
22 files changed, 311 insertions(+), 67 deletions(-)
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 174235d..a499755 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,11 @@
2014-11-13 Thomas Schwinge <tho...@codesourcery.com>
+ * gimplify.c (gimplify_omp_workshare) <OACC_DATA, OACC_KERNELS,
+ OACC_PARALLEL>: Don't request ORT_TARGET_MAP_FORCE.
+ (enum gimplify_omp_var_data, enum omp_region_type): Remove
+ GOVD_MAP_FORCE, and ORT_TARGET_MAP_FORCE, respectively. Update
+ all users.
+
* omp-low.c (scan_sharing_clauses) <OMP_CLAUSE_MAP>: Revert
earlier change.
diff --git gcc/gimplify.c gcc/gimplify.c
index 233ac56..2c8c666 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -94,8 +94,6 @@ enum gimplify_omp_var_data
/* Flags for GOVD_MAP. */
/* Don't copy back. */
GOVD_MAP_TO_ONLY = 8192,
- /* Force a specific behavior (or else, a run-time error). */
- GOVD_MAP_FORCE = 16384,
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
@@ -116,9 +114,7 @@ enum omp_region_type
/* Flags for ORT_TARGET. */
/* Prepare this region for offloading. */
- ORT_TARGET_OFFLOAD = 32,
- /* Default to GOVD_MAP_FORCE for implicit mappings in this region. */
- ORT_TARGET_MAP_FORCE = 64
+ ORT_TARGET_OFFLOAD = 32
};
/* Gimplify hashtable helper. */
@@ -5585,15 +5581,7 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree
decl, unsigned int flags)
if (!(flags & GOVD_LOCAL))
{
if (flags & GOVD_MAP)
- {
- nflags = GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT;
-#if 0
- /* Not sure if this is actually needed; haven't found a case
- where this would change anything; TODO. */
- if (flags & GOVD_MAP_FORCE)
- nflags |= OMP_CLAUSE_MAP_FORCE;
-#endif
- }
+ nflags = GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT;
else if (flags & GOVD_PRIVATE)
nflags = GOVD_PRIVATE;
else
@@ -5667,8 +5655,6 @@ omp_notice_threadprivate_variable (struct
gimplify_omp_ctx *ctx, tree decl,
if ((octx->region_type & ORT_TARGET)
&& (octx->region_type & ORT_TARGET_OFFLOAD))
{
- gcc_assert (!(octx->region_type & ORT_TARGET_MAP_FORCE));
-
n = splay_tree_lookup (octx->variables, (splay_tree_key)decl);
if (n == NULL)
{
@@ -5731,11 +5717,6 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree
decl, bool in_code)
if ((ctx->region_type & ORT_TARGET)
&& (ctx->region_type & ORT_TARGET_OFFLOAD))
{
- unsigned map_force;
- if (ctx->region_type & ORT_TARGET_MAP_FORCE)
- map_force = GOVD_MAP_FORCE;
- else
- map_force = 0;
ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
if (n == NULL)
{
@@ -5743,32 +5724,13 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree
decl, bool in_code)
{
error ("%qD referenced in target region does not have "
"a mappable type", decl);
- omp_add_variable (ctx, decl, GOVD_MAP | map_force | GOVD_EXPLICIT
| flags);
+ omp_add_variable (ctx, decl, GOVD_MAP | GOVD_EXPLICIT | flags);
}
else
- omp_add_variable (ctx, decl, GOVD_MAP | map_force | flags);
+ omp_add_variable (ctx, decl, GOVD_MAP | flags);
}
else
{
-#if 0
- /* The following fails for:
-
- int l = 10;
- float c[l];
- #pragma acc parallel copy(c[2:4])
- {
- #pragma acc parallel
- {
- int t = sizeof c;
- }
- }
-
- ..., which we currently don't have to care about (nesting
- disabled), but eventually will have to; TODO. */
- if ((n->value & GOVD_MAP) && !(n->value & GOVD_EXPLICIT))
- gcc_assert ((n->value & GOVD_MAP_FORCE) == map_force);
-#endif
-
/* If nothing changed, there's nothing left to do. */
if ((n->value & flags) == flags)
return ret;
@@ -6423,13 +6385,11 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void
*data)
OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1;
else if (code == OMP_CLAUSE_MAP)
{
- unsigned map_kind;
+ enum omp_clause_map_kind map_kind;
map_kind = (flags & GOVD_MAP_TO_ONLY
? OMP_CLAUSE_MAP_TO
: OMP_CLAUSE_MAP_TOFROM);
- if (flags & GOVD_MAP_FORCE)
- map_kind |= OMP_CLAUSE_MAP_FORCE;
- OMP_CLAUSE_MAP_KIND (clause) = (enum omp_clause_map_kind) map_kind;
+ OMP_CLAUSE_MAP_KIND (clause) = map_kind;
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
@@ -7258,23 +7218,16 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
switch (TREE_CODE (expr))
{
- case OACC_DATA:
- ort = (enum omp_region_type) (ORT_TARGET
- | ORT_TARGET_MAP_FORCE);
- break;
- case OACC_KERNELS:
- case OACC_PARALLEL:
- ort = (enum omp_region_type) (ORT_TARGET
- | ORT_TARGET_OFFLOAD
- | ORT_TARGET_MAP_FORCE);
- break;
case OMP_SECTIONS:
case OMP_SINGLE:
ort = ORT_WORKSHARE;
break;
+ case OACC_KERNELS:
+ case OACC_PARALLEL:
case OMP_TARGET:
ort = (enum omp_region_type) (ORT_TARGET | ORT_TARGET_OFFLOAD);
break;
+ case OACC_DATA:
case OMP_TARGET_DATA:
ort = ORT_TARGET;
break;
diff --git include/ChangeLog.gomp include/ChangeLog.gomp
new file mode 100644
index 0000000..9172c26
--- /dev/null
+++ include/ChangeLog.gomp
@@ -0,0 +1,4 @@
+2014-11-13 Thomas Schwinge <tho...@codesourcery.com>
+
+ * gomp-constants.h: Define _GOMP_MAP_FLAG_SPECIAL and
+ _GOMP_MAP_FLAG_FORCE.
diff --git include/gomp-constants.h include/gomp-constants.h
index e600766..15b658f 100644
--- include/gomp-constants.h
+++ include/gomp-constants.h
@@ -28,6 +28,9 @@
/* Enumerated variable mapping types used to communicate between GCC and
libgomp. These values are used for both OpenMP and OpenACC. */
+#define _GOMP_MAP_FLAG_SPECIAL (1 << 2)
+#define _GOMP_MAP_FLAG_FORCE (1 << 3)
+
#define GOMP_MAP_ALLOC 0x00
#define GOMP_MAP_ALLOC_TO 0x01
#define GOMP_MAP_ALLOC_FROM 0x02
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 0528531..254846f 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,3 +1,26 @@
+2014-11-13 Thomas Schwinge <tho...@codesourcery.com>
+
+ * target.c (gomp_map_vars_existing): Error out if "force"
+ semantics.
+ (gomp_map_vars): Actually pass kinds to gomp_map_vars_existing.
+ Remove FIXMEs.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-1.c: New file.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-3.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-4.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-5.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-6.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-7.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-1.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-2.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-3.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-4.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-5.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-6.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-7.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-8.f: Likewise.
+
2014-11-12 Thomas Schwinge <tho...@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/collapse-4.c: New file.
diff --git libgomp/target.c libgomp/target.c
index 052c59d..2b9f08f 100644
--- libgomp/target.c
+++ libgomp/target.c
@@ -117,9 +117,11 @@ static inline void
gomp_map_vars_existing (splay_tree_key oldn, splay_tree_key newn,
unsigned char kind)
{
- if (oldn->host_start > newn->host_start
+ if ((!(kind & _GOMP_MAP_FLAG_SPECIAL)
+ && (kind & _GOMP_MAP_FLAG_FORCE))
+ || oldn->host_start > newn->host_start
|| oldn->host_end < newn->host_end)
- gomp_fatal ("Trying to map into device [%p..%p) object when"
+ gomp_fatal ("Trying to map into device [%p..%p) object when "
"[%p..%p) is already mapped",
(void *) newn->host_start, (void *) newn->host_end,
(void *) oldn->host_start, (void *) oldn->host_end);
@@ -200,7 +202,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t
mapnum,
if (n)
{
tgt->list[i] = n;
- gomp_map_vars_existing (n, &cur_node, kind);
+ gomp_map_vars_existing (n, &cur_node, kind & typemask);
}
else
{
@@ -323,7 +325,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t
mapnum,
if (n)
{
tgt->list[i] = n;
- gomp_map_vars_existing (n, k, kind);
+ gomp_map_vars_existing (n, k, kind & typemask);
}
else
{
@@ -345,18 +347,15 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t
mapnum,
switch (kind & typemask)
{
- case GOMP_MAP_FORCE_ALLOC:
- case GOMP_MAP_FORCE_FROM:
- /* FIXME: No special handling (see comment in
- oacc-parallel.c). */
case GOMP_MAP_ALLOC:
case GOMP_MAP_ALLOC_FROM:
+ case GOMP_MAP_FORCE_ALLOC:
+ case GOMP_MAP_FORCE_FROM:
break;
- case GOMP_MAP_FORCE_TO:
- case GOMP_MAP_FORCE_TOFROM:
- /* FIXME: No special handling, as above. */
case GOMP_MAP_ALLOC_TO:
case GOMP_MAP_ALLOC_TOFROM:
+ case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_FORCE_TOFROM:
/* Copy from host to device memory. */
/* FIXME: Perhaps add some smarts, like if copying
several adjacent fields from host to target, use some
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c
libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c
new file mode 100644
index 0000000..83c0a42
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c
@@ -0,0 +1,19 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <openacc.h>
+
+int
+main (int argc, char *argv[])
+{
+ int i;
+
+ acc_copyin (&i, sizeof i);
+
+#pragma acc data copy (i)
+ ++i;
+
+ return 0;
+}
+
+/* { dg-shouldfail "" }
+ { dg-output "Trying to map into device .* object when .* is already mapped"
} */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c
libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c
new file mode 100644
index 0000000..137d8ce
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c
@@ -0,0 +1,16 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+int
+main (int argc, char *argv[])
+{
+ int i;
+
+#pragma acc data present_or_copy (i)
+#pragma acc data copyout (i)
+ ++i;
+
+ return 0;
+}
+
+/* { dg-shouldfail "" }
+ { dg-output "Trying to map into device .* object when .* is already mapped"
} */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c
libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c
new file mode 100644
index 0000000..b993b78
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c
@@ -0,0 +1,17 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <openacc.h>
+
+int
+main (int argc, char *argv[])
+{
+ int i;
+
+#pragma acc data present_or_copy (i)
+ acc_copyin (&i, sizeof i);
+
+ return 0;
+}
+
+/* { dg-shouldfail "" }
+ { dg-output "already mapped to" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c
libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c
new file mode 100644
index 0000000..82523f4
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c
@@ -0,0 +1,17 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <openacc.h>
+
+int
+main (int argc, char *argv[])
+{
+ int i;
+
+ acc_present_or_copyin (&i, sizeof i);
+ acc_copyin (&i, sizeof i);
+
+ return 0;
+}
+
+/* { dg-shouldfail "" }
+ { dg-output "already mapped to" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c
libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c
new file mode 100644
index 0000000..4961fe5
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c
@@ -0,0 +1,17 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <openacc.h>
+
+int
+main (int argc, char *argv[])
+{
+ int i;
+
+#pragma acc enter data create (i)
+ acc_copyin (&i, sizeof i);
+
+ return 0;
+}
+
+/* { dg-shouldfail "" }
+ { dg-output "already mapped to" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c
libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c
new file mode 100644
index 0000000..77b56a9
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c
@@ -0,0 +1,17 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <openacc.h>
+
+int
+main (int argc, char *argv[])
+{
+ int i;
+
+ acc_present_or_copyin (&i, sizeof i);
+#pragma acc enter data create (i)
+
+ return 0;
+}
+
+/* { dg-shouldfail "" }
+ { dg-output "already mapped to" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c
libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c
new file mode 100644
index 0000000..b08417b
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c
@@ -0,0 +1,17 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <openacc.h>
+
+int
+main (int argc, char *argv[])
+{
+ int i;
+
+#pragma acc enter data create (i)
+ acc_create (&i, sizeof i);
+
+ return 0;
+}
+
+/* { dg-shouldfail "" }
+ { dg-output "already mapped to" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c
libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c
new file mode 100644
index 0000000..a50f7de
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c
@@ -0,0 +1,16 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+int
+main (int argc, char *argv[])
+{
+ int i;
+
+#pragma acc data create (i)
+#pragma acc parallel copyin (i)
+ ++i;
+
+ return 0;
+}
+
+/* { dg-shouldfail "" }
+ { dg-output "Trying to map into device .* object when .* is already mapped"
} */
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f
libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f
new file mode 100644
index 0000000..ac220ab
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f
@@ -0,0 +1,17 @@
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+ IMPLICIT NONE
+ INCLUDE "openacc_lib.h"
+
+ INTEGER I
+
+ CALL ACC_COPYIN (I)
+
+!$ACC DATA COPY (I)
+ I = 0
+!$ACC END DATA
+
+ END
+
+! { dg-shouldfail "" }
+! { dg-output "Trying to map into device .* object when .* is already mapped" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f
libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f
new file mode 100644
index 0000000..2c5254b
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f
@@ -0,0 +1,16 @@
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+ IMPLICIT NONE
+
+ INTEGER I
+
+!$ACC DATA PRESENT_OR_COPY (I)
+!$ACC DATA COPYOUT (I)
+ I = 0
+!$ACC END DATA
+!$ACC END DATA
+
+ END
+
+! { dg-shouldfail "" }
+! { dg-output "Trying to map into device .* object when .* is already mapped" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f
libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f
new file mode 100644
index 0000000..c41de28
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f
@@ -0,0 +1,15 @@
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+ IMPLICIT NONE
+ INCLUDE "openacc_lib.h"
+
+ INTEGER I
+
+!$ACC DATA PRESENT_OR_COPY (I)
+ CALL ACC_COPYIN (I)
+!$ACC END DATA
+
+ END
+
+! { dg-shouldfail "" }
+! { dg-output "already mapped to" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f
libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f
new file mode 100644
index 0000000..f54bf58
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f
@@ -0,0 +1,14 @@
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+ IMPLICIT NONE
+ INCLUDE "openacc_lib.h"
+
+ INTEGER I
+
+ CALL ACC_PRESENT_OR_COPYIN (I)
+ CALL ACC_COPYIN (I)
+
+ END
+
+! { dg-shouldfail "" }
+! { dg-output "already mapped to" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f
libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f
new file mode 100644
index 0000000..9a3e94f
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f
@@ -0,0 +1,14 @@
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+ IMPLICIT NONE
+ INCLUDE "openacc_lib.h"
+
+ INTEGER I
+
+!$ACC ENTER DATA CREATE (I)
+ CALL ACC_COPYIN (I)
+
+ END
+
+! { dg-shouldfail "" }
+! { dg-output "already mapped to" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f
libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f
new file mode 100644
index 0000000..eaf5d98
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f
@@ -0,0 +1,14 @@
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+ IMPLICIT NONE
+ INCLUDE "openacc_lib.h"
+
+ INTEGER I
+
+ CALL ACC_PRESENT_OR_COPYIN (I)
+!$ACC ENTER DATA CREATE (I)
+
+ END
+
+! { dg-shouldfail "" }
+! { dg-output "already mapped to" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f
libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f
new file mode 100644
index 0000000..d96bf0b
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f
@@ -0,0 +1,14 @@
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+ IMPLICIT NONE
+ INCLUDE "openacc_lib.h"
+
+ INTEGER I
+
+!$ACC ENTER DATA CREATE (I)
+ CALL ACC_CREATE (I)
+
+ END
+
+! { dg-shouldfail "" }
+! { dg-output "already mapped to" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f
libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f
new file mode 100644
index 0000000..16da048
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f
@@ -0,0 +1,16 @@
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+ IMPLICIT NONE
+
+ INTEGER I
+
+!$ACC DATA CREATE (I)
+!$ACC PARALLEL COPYIN (I)
+ I = 0
+!$ACC END PARALLEL
+!$ACC END DATA
+
+ END
+
+! { dg-shouldfail "" }
+! { dg-output "Trying to map into device .* object when .* is already mapped" }
Grüße,
Thomas
signature.asc
Description: PGP signature
