diff mbox

[gomp4,3/6] Initial support for OpenACC memory mapping semantics.

Message ID 87fvdncmxg.fsf@kepler.schwinge.homeip.net
State New
Headers show

Commit Message

Thomas Schwinge Nov. 13, 2014, 12:19 p.m. UTC 
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(-)



Grüße,
 Thomas

Comments

Jakub Jelinek Nov. 13, 2014, 1:10 p.m. UTC | #1 
On Thu, Nov 13, 2014 at 01:19:55PM +0100, Thomas Schwinge wrote:
> --- 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)

I'm worried about reserved namespace issues if you use _ followed by
capital letter.  Can't it be just GOMP_MAP_FLAG_* ?

	Jakub
Thomas Schwinge Nov. 13, 2014, 1:38 p.m. UTC | #2 
Hi Jakub!

On Thu, 13 Nov 2014 14:10:10 +0100, Jakub Jelinek <jakub@redhat.com> wrote:
> On Thu, Nov 13, 2014 at 01:19:55PM +0100, Thomas Schwinge wrote:
> > --- 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)
> 
> I'm worried about reserved namespace issues if you use _ followed by
> capital letter.

Please remind me what those are reserved for?

>  Can't it be just GOMP_MAP_FLAG_* ?

My worry is the other way round: gomp-constants.h is also #included from
<openacc.h> (to grab some of its constants), and using plain GOMP_* would
pollute the user's namespace?  (I'm working on a patch to clean that up,
and also use gomp-constants.h more often, also for OpenMP code.)  (Such a
shared (GCC/libgomp) header files had been discussed before, and now
introduced in
<http://news.gmane.org/find-root.php?message_id=%3C20140923191931.2177e60f%40octopus%3E>.)


Grüße,
 Thomas
Jakub Jelinek Nov. 13, 2014, 1:51 p.m. UTC | #3 
On Thu, Nov 13, 2014 at 02:38:06PM +0100, Thomas Schwinge wrote:
> Hi Jakub!
> 
> On Thu, 13 Nov 2014 14:10:10 +0100, Jakub Jelinek <jakub@redhat.com> wrote:
> > On Thu, Nov 13, 2014 at 01:19:55PM +0100, Thomas Schwinge wrote:
> > > --- 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)
> > 
> > I'm worried about reserved namespace issues if you use _ followed by
> > capital letter.
> 
> Please remind me what those are reserved for?

See e.g.
http://www.gnu.org/software/libc/manual/html_node/Reserved-Names.html
http://pubs.opengroup.org/onlinepubs/007904975/functions/xsh_chap02_02.html
and remember that if you use gomp-constants.h in the compiler, it can be
built by the system compiler, which can be a very different implementation.

> >  Can't it be just GOMP_MAP_FLAG_* ?
> 
> My worry is the other way round: gomp-constants.h is also #included from
> <openacc.h> (to grab some of its constants), and using plain GOMP_* would
> pollute the user's namespace?  (I'm working on a patch to clean that up,
> and also use gomp-constants.h more often, also for OpenMP code.)  (Such a
> shared (GCC/libgomp) header files had been discussed before, and now
> introduced in
> <http://news.gmane.org/find-root.php?message_id=%3C20140923191931.2177e60f%40octopus%3E>.)

I think including gomp-constants.h in openacc.h, if that is a publicly
installed header, is a bad idea, you'll pollute namespace of that header.
Just duplicate the values in there under the right standard required names,
and you want, either add a testcase or some static assertions (e.g.
of the kind extern char typedef1[condition ? 1 : -1]; in some macros)
to verify that the openacc.h constants match the gomp-constants.h where
required.

	Jakub
diff mbox

Patch

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  <thomas@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  <thomas@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  <thomas@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  <thomas@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" }