diff mbox

[gomp4.1] map clause parsing improvements

Message ID 20150611121420.GY10247@tucnak.redhat.com
State New
Headers show

Commit Message

Jakub Jelinek June 11, 2015, 12:14 p.m. UTC 
On Tue, Jun 09, 2015 at 09:36:08PM +0300, Ilya Verbin wrote:
> On Wed, Apr 29, 2015 at 14:06:44 +0200, Jakub Jelinek wrote:
> > [...] The draft requires only alloc or to
> > (or always, variants) for enter data and only from or delete (or always,
> > variants) for exit data, so in theory it is possible to figure that from
> > the call without extra args, but not so for update - enter data is supposed
> > to increment reference counts, exit data decrement. [...]
> 
> TR3.pdf also says about 'release' map-type for exit data, but it is not
> described in the document.

So, I've committed a patch to add parsing release map-kind, and fix up or add
verification in C/C++ FE what map-kinds are used.

Furthermore, it seems the OpenMP 4.1 always modifier is something completely
unrelated to the OpenACC force flag, in OpenMP 4.1 everything is reference
count based, and always seems to make a difference only for from/to/tofrom,
where it says that the copying is done unconditionally; thus the patch uses
a different bit for that.

For array sections resulting in GOMP_MAP_POINTER kind perhaps we'll also
need an always variant for it, and supposedly on the exit data construct
we'll want to turn GOMP_MAP_POINTER into GOMP_MAP_RELEASE or GOMP_MAP_DELETE
depending on what it was originally.

2015-06-11  Jakub Jelinek  <jakub@redhat.com>

include/
	* gomp-constants.h (GOMP_MAP_FLAG_ALWAYS): Define.
	(enum gomp_map_kind): Add GOMP_MAP_ALWAYS_TO, GOMP_MAP_ALWAYS_FROM,
	GOMP_MAP_ALWAYS_TOFROM, GOMP_MAP_DELETE, GOMP_MAP_RELEASE.
gcc/
	* omp-low.c (lower_omp_target): Accept GOMP_MAP_RELEASE,
	GOMP_MAP_ALWAYS_TO, GOMP_MAP_ALWAYS_FROM and GOMP_MAP_ALWAYS_TOFROM.
	Accept GOMP_MAP_FORCE* except for FORCE_DEALLOC only for OpenACC.
	* tree-pretty-print.c (dump_omp_clause): Print GOMP_MAP_FORCE_*
	as force_*, handle GOMP_MAP_ALWAYS_*.
gcc/c/
	* c-parser.c (c_parser_omp_clause_map): Handle release
	map kind.  Adjust to use GOMP_MAP_ALWAYS_* and only on
	clauses where it makes sense.
	(c_parser_omp_target_data): Diagnose invalid map-kind
	for the construct.
	(c_parser_omp_target_enter_data): Adjust for new encoding
	of always,from, always,to and always,tofrom.  Accept
	GOMP_MAP_POINTER.
	(c_parser_omp_target_exit_data): Likewise.
	(c_parser_omp_target): Diagnose invalid map-kind for the
	construct.
gcc/cp/
	* parser.c (cp_parser_omp_clause_map): Handle release
	map kind.  Adjust to use GOMP_MAP_ALWAYS_* and only on
	clauses where it makes sense.
	(cp_parser_omp_target_data): Diagnose invalid map-kind
	for the construct.
	(cp_parser_omp_target_enter_data): Adjust for new encoding
	of always,from, always,to and always,tofrom.  Accept
	GOMP_MAP_POINTER.
	(cp_parser_omp_target_exit_data): Likewise.
	(cp_parser_omp_target): Diagnose invalid map-kind for the
	construct.



	Jakub

Comments

Thomas Schwinge Oct. 19, 2015, 10:20 a.m. UTC | #1 
Hi!

On Thu, 11 Jun 2015 14:14:20 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
> On Tue, Jun 09, 2015 at 09:36:08PM +0300, Ilya Verbin wrote:
> > On Wed, Apr 29, 2015 at 14:06:44 +0200, Jakub Jelinek wrote:
> > > [...] The draft requires only alloc or to
> > > (or always, variants) for enter data and only from or delete (or always,
> > > variants) for exit data, so in theory it is possible to figure that from
> > > the call without extra args, but not so for update - enter data is supposed
> > > to increment reference counts, exit data decrement. [...]
> > 
> > TR3.pdf also says about 'release' map-type for exit data, but it is not
> > described in the document.
> 
> So, I've committed a patch to add parsing release map-kind, and fix up or add
> verification in C/C++ FE what map-kinds are used.
> 
> Furthermore, it seems the OpenMP 4.1 always modifier is something completely
> unrelated to the OpenACC force flag, in OpenMP 4.1 everything is reference
> count based, and always seems to make a difference only for from/to/tofrom,
> where it says that the copying is done unconditionally; thus the patch uses
> a different bit for that.

Aha, I see.  (The poor OpenACC/OpenMP users, having to remember so may
small yet intricate details...)

> include/
> 	* gomp-constants.h (GOMP_MAP_FLAG_ALWAYS): Define.
> 	(enum gomp_map_kind): Add GOMP_MAP_ALWAYS_TO, GOMP_MAP_ALWAYS_FROM,
> 	GOMP_MAP_ALWAYS_TOFROM, GOMP_MAP_DELETE, GOMP_MAP_RELEASE.

> --- include/gomp-constants.h.jj	2015-05-21 11:12:09.000000000 +0200
> +++ include/gomp-constants.h	2015-06-11 11:24:32.041654947 +0200
> @@ -41,6 +41,8 @@
>  #define GOMP_MAP_FLAG_SPECIAL_1		(1 << 3)
>  #define GOMP_MAP_FLAG_SPECIAL		(GOMP_MAP_FLAG_SPECIAL_1 \
>  					 | GOMP_MAP_FLAG_SPECIAL_0)
> +/* OpenMP always flag.  */
> +#define GOMP_MAP_FLAG_ALWAYS		(1 << 6)
>  /* Flag to force a specific behavior (or else, trigger a run-time error).  */
>  #define GOMP_MAP_FLAG_FORCE		(1 << 7)
>  
> @@ -77,7 +79,21 @@ enum gomp_map_kind
>      /* ..., and copy from device.  */
>      GOMP_MAP_FORCE_FROM =		(GOMP_MAP_FLAG_FORCE | GOMP_MAP_FROM),
>      /* ..., and copy to and from device.  */
> -    GOMP_MAP_FORCE_TOFROM =		(GOMP_MAP_FLAG_FORCE | GOMP_MAP_TOFROM)
> +    GOMP_MAP_FORCE_TOFROM =		(GOMP_MAP_FLAG_FORCE | GOMP_MAP_TOFROM),
> +    /* If not already present, allocate.  And unconditionally copy to
> +       device.  */
> +    GOMP_MAP_ALWAYS_TO =		(GOMP_MAP_FLAG_ALWAYS | GOMP_MAP_TO),
> +    /* If not already present, allocate.  And unconditionally copy from
> +       device.  */
> +    GOMP_MAP_ALWAYS_FROM =		(GOMP_MAP_FLAG_ALWAYS | GOMP_MAP_FROM),
> +    /* If not already present, allocate.  And unconditionally copy to and from
> +       device.  */
> +    GOMP_MAP_ALWAYS_TOFROM =		(GOMP_MAP_FLAG_ALWAYS | GOMP_MAP_TOFROM),
> +    /* OpenMP 4.1 alias for forced deallocation.  */
> +    GOMP_MAP_DELETE =			GOMP_MAP_FORCE_DEALLOC,

To avoid confusion about two different identifiers naming the same
functionality, I'd prefer to avoid such aliases ("GOMP_MAP_DELETE =
GOMP_MAP_FORCE_DEALLOC"), and instead just rename GOMP_MAP_FORCE_DEALLOC
to GOMP_MAP_DELETE, if that's the name you prefer.

By the way, looking at GCC 6 libgomp compatibility regarding
OpenACC/nvptx offloading for executables compiled with GCC 5, for the
legacy entry point libgomp/oacc-parallel.c:GOACC_parallel only supports
host-fallback execution, which doesn't pay attention to data clause at
all (sizes and kinds formal parameters), so you're free to renumber
GOMP_MAP_* if/where that makes sense.

> +    /* Decrement usage count and deallocate if zero.  */
> +    GOMP_MAP_RELEASE =			(GOMP_MAP_FLAG_ALWAYS
> +					 | GOMP_MAP_FORCE_DEALLOC)
>    };

I have not yet read the OpenMP 4.1/4.5 standard, but it's not obvious to
me here how the GOMP_MAP_FLAG_ALWAYS flag relates to the OpenMP release
clause (GOMP_MAP_RELEASE here)?  Shouldn't GOMP_MAP_RELEASE be
"(GOMP_MAP_FLAG_SPECIAL_1 | 3)" or similar?


Grüße
 Thomas
Jakub Jelinek Oct. 19, 2015, 10:34 a.m. UTC | #2 
On Mon, Oct 19, 2015 at 12:20:23PM +0200, Thomas Schwinge wrote:
> > @@ -77,7 +79,21 @@ enum gomp_map_kind
> >      /* ..., and copy from device.  */
> >      GOMP_MAP_FORCE_FROM =		(GOMP_MAP_FLAG_FORCE | GOMP_MAP_FROM),
> >      /* ..., and copy to and from device.  */
> > -    GOMP_MAP_FORCE_TOFROM =		(GOMP_MAP_FLAG_FORCE | GOMP_MAP_TOFROM)
> > +    GOMP_MAP_FORCE_TOFROM =		(GOMP_MAP_FLAG_FORCE | GOMP_MAP_TOFROM),
> > +    /* If not already present, allocate.  And unconditionally copy to
> > +       device.  */
> > +    GOMP_MAP_ALWAYS_TO =		(GOMP_MAP_FLAG_ALWAYS | GOMP_MAP_TO),
> > +    /* If not already present, allocate.  And unconditionally copy from
> > +       device.  */
> > +    GOMP_MAP_ALWAYS_FROM =		(GOMP_MAP_FLAG_ALWAYS | GOMP_MAP_FROM),
> > +    /* If not already present, allocate.  And unconditionally copy to and from
> > +       device.  */
> > +    GOMP_MAP_ALWAYS_TOFROM =		(GOMP_MAP_FLAG_ALWAYS | GOMP_MAP_TOFROM),
> > +    /* OpenMP 4.1 alias for forced deallocation.  */
> > +    GOMP_MAP_DELETE =			GOMP_MAP_FORCE_DEALLOC,
> 
> To avoid confusion about two different identifiers naming the same
> functionality, I'd prefer to avoid such aliases ("GOMP_MAP_DELETE =
> GOMP_MAP_FORCE_DEALLOC"), and instead just rename GOMP_MAP_FORCE_DEALLOC
> to GOMP_MAP_DELETE, if that's the name you prefer.

If you are ok with removing GOMP_MAP_FORCE_DEALLOC and just use
GOMP_MAP_DELETE, that is ok by me, just post a patch.

> By the way, looking at GCC 6 libgomp compatibility regarding
> OpenACC/nvptx offloading for executables compiled with GCC 5, for the
> legacy entry point libgomp/oacc-parallel.c:GOACC_parallel only supports
> host-fallback execution, which doesn't pay attention to data clause at
> all (sizes and kinds formal parameters), so you're free to renumber
> GOMP_MAP_* if/where that makes sense.
> 
> > +    /* Decrement usage count and deallocate if zero.  */
> > +    GOMP_MAP_RELEASE =			(GOMP_MAP_FLAG_ALWAYS
> > +					 | GOMP_MAP_FORCE_DEALLOC)
> >    };
> 
> I have not yet read the OpenMP 4.1/4.5 standard, but it's not obvious to
> me here how the GOMP_MAP_FLAG_ALWAYS flag relates to the OpenMP release
> clause (GOMP_MAP_RELEASE here)?  Shouldn't GOMP_MAP_RELEASE be
> "(GOMP_MAP_FLAG_SPECIAL_1 | 3)" or similar?

It isn't related to always, but always really is something that affects
solely the data movement (i.e. to, from, tofrom), and while it can be
specified elsewhere, it makes no difference.  Wasting one bit just for that
is something we don't have the luxury for, which is why I've started using
that bit for other OpenMP stuff (it acts there like GOMP_MAP_FLAG_SPECIAL_2
to some extent).  It is not just release, but also the struct mapping etc.
I'll still need to make further changes, because the rules for mapping
structure element pointer/reference based array sections and structure
element references have changed again.

Some changes in the enum can be of course still be done until say mid stage3
but at least for OpenMP 4.0 we should keep backwards compatibility (so
whatever we've already used in GCC 4.9/5 should keep working).

	Jakub
Thomas Schwinge Oct. 19, 2015, 3 p.m. UTC | #3 
Hi!

On Mon, 19 Oct 2015 12:34:08 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
> On Mon, Oct 19, 2015 at 12:20:23PM +0200, Thomas Schwinge wrote:
> > > +    /* Decrement usage count and deallocate if zero.  */
> > > +    GOMP_MAP_RELEASE =			(GOMP_MAP_FLAG_ALWAYS
> > > +					 | GOMP_MAP_FORCE_DEALLOC)
> > >    };
> > 
> > I have not yet read the OpenMP 4.1/4.5 standard, but it's not obvious to
> > me here how the GOMP_MAP_FLAG_ALWAYS flag relates to the OpenMP release
> > clause (GOMP_MAP_RELEASE here)?  Shouldn't GOMP_MAP_RELEASE be
> > "(GOMP_MAP_FLAG_SPECIAL_1 | 3)" or similar?
> 
> It isn't related to always, but always really is something that affects
> solely the data movement (i.e. to, from, tofrom), and while it can be
> specified elsewhere, it makes no difference.  Wasting one bit just for that
> is something we don't have the luxury for, which is why I've started using
> that bit for other OpenMP stuff (it acts there like GOMP_MAP_FLAG_SPECIAL_2
> to some extent).  It is not just release, but also the struct mapping etc.
> I'll still need to make further changes, because the rules for mapping
> structure element pointer/reference based array sections and structure
> element references have changed again.

Hmm, I do think we should allow the luxury to use its own bit for
GOMP_MAP_FLAG_ALWAYS -- we can extend the interface later, should we
really find uses for the other two remaining bits -- or if not using a
separate bit, at least make sure that GOMP_MAP_FLAG_ALWAYS is not used as
a flag.  See, for example, the following occasions where
GOMP_MAP_FLAG_ALWAYS is used as a flag: these conditionals will also be
matched for GOMP_MAP_STRUCT, GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION, and
GOMP_MAP_RELEASE.  I have not analyzed whether that is erroneous or not,
but it surely is confusing?

    $ < gcc/gimplify.c grep -C3 GOMP_MAP_FLAG_ALWAYS
                          struct_map_to_clause->put (decl, *list_p);
                          list_p = &OMP_CLAUSE_CHAIN (*list_p);
                          flags = GOVD_MAP | GOVD_EXPLICIT;
                          if (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS)
                            flags |= GOVD_SEEN;
                          goto do_add_decl;
                        }
    --
                          tree *sc = NULL, *pt = NULL;
                          if (!ptr && TREE_CODE (*osc) == TREE_LIST)
                            osc = &TREE_PURPOSE (*osc);
                          if (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS)
                            n->value |= GOVD_SEEN;
                          offset_int o1, o2;
                          if (offset)
    --
              n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
              if ((ctx->region_type & ORT_TARGET) != 0
                  && !(n->value & GOVD_SEEN)
                  && ((OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS) == 0
                      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT))
                {
                  remove = true;

I'd suggest turning GOMP_MAP_FLAG_ALWAYS into GOMP_MAP_FLAG_SPECIAL_2,
and then provide a GOMP_MAP_ALWAYS_P that evaluates to true just for the
three "always,to", "always,from", and "always,tofrom" cases.


Grüße,
 Thomas
Jakub Jelinek Oct. 20, 2015, 10:03 a.m. UTC | #4 
On Mon, Oct 19, 2015 at 05:00:33PM +0200, Thomas Schwinge wrote:
>               n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
>               if ((ctx->region_type & ORT_TARGET) != 0
>                   && !(n->value & GOVD_SEEN)
>                   && ((OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS) == 0
>                       || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT))
>                 {
>                   remove = true;
> 
> I'd suggest turning GOMP_MAP_FLAG_ALWAYS into GOMP_MAP_FLAG_SPECIAL_2,
> and then provide a GOMP_MAP_ALWAYS_P that evaluates to true just for the
> three "always,to", "always,from", and "always,tofrom" cases.

Yeah, that can be done, I'll add it to my todo list.

	Jakub
diff mbox

Patch

--- include/gomp-constants.h.jj	2015-05-21 11:12:09.000000000 +0200
+++ include/gomp-constants.h	2015-06-11 11:24:32.041654947 +0200
@@ -41,6 +41,8 @@ 
 #define GOMP_MAP_FLAG_SPECIAL_1		(1 << 3)
 #define GOMP_MAP_FLAG_SPECIAL		(GOMP_MAP_FLAG_SPECIAL_1 \
 					 | GOMP_MAP_FLAG_SPECIAL_0)
+/* OpenMP always flag.  */
+#define GOMP_MAP_FLAG_ALWAYS		(1 << 6)
 /* Flag to force a specific behavior (or else, trigger a run-time error).  */
 #define GOMP_MAP_FLAG_FORCE		(1 << 7)
 
@@ -77,7 +79,21 @@  enum gomp_map_kind
     /* ..., and copy from device.  */
     GOMP_MAP_FORCE_FROM =		(GOMP_MAP_FLAG_FORCE | GOMP_MAP_FROM),
     /* ..., and copy to and from device.  */
-    GOMP_MAP_FORCE_TOFROM =		(GOMP_MAP_FLAG_FORCE | GOMP_MAP_TOFROM)
+    GOMP_MAP_FORCE_TOFROM =		(GOMP_MAP_FLAG_FORCE | GOMP_MAP_TOFROM),
+    /* If not already present, allocate.  And unconditionally copy to
+       device.  */
+    GOMP_MAP_ALWAYS_TO =		(GOMP_MAP_FLAG_ALWAYS | GOMP_MAP_TO),
+    /* If not already present, allocate.  And unconditionally copy from
+       device.  */
+    GOMP_MAP_ALWAYS_FROM =		(GOMP_MAP_FLAG_ALWAYS | GOMP_MAP_FROM),
+    /* If not already present, allocate.  And unconditionally copy to and from
+       device.  */
+    GOMP_MAP_ALWAYS_TOFROM =		(GOMP_MAP_FLAG_ALWAYS | GOMP_MAP_TOFROM),
+    /* OpenMP 4.1 alias for forced deallocation.  */
+    GOMP_MAP_DELETE =			GOMP_MAP_FORCE_DEALLOC,
+    /* Decrement usage count and deallocate if zero.  */
+    GOMP_MAP_RELEASE =			(GOMP_MAP_FLAG_ALWAYS
+					 | GOMP_MAP_FORCE_DEALLOC)
   };
 
 #define GOMP_MAP_COPY_TO_P(X) \
--- gcc/omp-low.c.jj	2015-06-10 19:50:26.000000000 +0200
+++ gcc/omp-low.c	2015-06-11 11:35:02.515892363 +0200
@@ -12511,13 +12511,17 @@  lower_omp_target (gimple_stmt_iterator *
 	  case GOMP_MAP_TOFROM:
 	  case GOMP_MAP_POINTER:
 	  case GOMP_MAP_TO_PSET:
+	  case GOMP_MAP_FORCE_DEALLOC:
+	  case GOMP_MAP_RELEASE:
+	  case GOMP_MAP_ALWAYS_TO:
+	  case GOMP_MAP_ALWAYS_FROM:
+	  case GOMP_MAP_ALWAYS_TOFROM:
+	    break;
 	  case GOMP_MAP_FORCE_ALLOC:
 	  case GOMP_MAP_FORCE_TO:
 	  case GOMP_MAP_FORCE_FROM:
 	  case GOMP_MAP_FORCE_TOFROM:
 	  case GOMP_MAP_FORCE_PRESENT:
-	  case GOMP_MAP_FORCE_DEALLOC:
-	    break;
 	  case GOMP_MAP_FORCE_DEVICEPTR:
 	    gcc_assert (is_gimple_omp_oacc (stmt));
 	    break;
--- gcc/tree-pretty-print.c.jj	2015-05-19 18:56:43.000000000 +0200
+++ gcc/tree-pretty-print.c	2015-06-11 11:32:39.814102121 +0200
@@ -560,26 +560,38 @@  dump_omp_clause (pretty_printer *pp, tre
 	  pp_string (pp, "tofrom");
 	  break;
 	case GOMP_MAP_FORCE_ALLOC:
-	  pp_string (pp, "always,alloc");
+	  pp_string (pp, "force_alloc");
 	  break;
 	case GOMP_MAP_FORCE_TO:
-	  pp_string (pp, "always,to");
+	  pp_string (pp, "force_to");
 	  break;
 	case GOMP_MAP_FORCE_FROM:
-	  pp_string (pp, "always,from");
+	  pp_string (pp, "force_from");
 	  break;
 	case GOMP_MAP_FORCE_TOFROM:
-	  pp_string (pp, "always,tofrom");
+	  pp_string (pp, "force_tofrom");
 	  break;
 	case GOMP_MAP_FORCE_PRESENT:
 	  pp_string (pp, "force_present");
 	  break;
 	case GOMP_MAP_FORCE_DEALLOC:
-	  pp_string (pp, "always,delete");
+	  pp_string (pp, "delete");
 	  break;
 	case GOMP_MAP_FORCE_DEVICEPTR:
 	  pp_string (pp, "force_deviceptr");
 	  break;
+	case GOMP_MAP_ALWAYS_TO:
+	  pp_string (pp, "always,to");
+	  break;
+	case GOMP_MAP_ALWAYS_FROM:
+	  pp_string (pp, "always,from");
+	  break;
+	case GOMP_MAP_ALWAYS_TOFROM:
+	  pp_string (pp, "always,tofrom");
+	  break;
+	case GOMP_MAP_RELEASE:
+	  pp_string (pp, "release");
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
--- gcc/c/c-parser.c.jj	2015-06-10 14:52:06.000000000 +0200
+++ gcc/c/c-parser.c	2015-06-11 13:02:30.788629315 +0200
@@ -11661,7 +11661,7 @@  c_parser_omp_clause_depend (c_parser *pa
 
    OpenMP 4.1:
    map-kind:
-     alloc | to | from | tofrom | delete
+     alloc | to | from | tofrom | release | delete
 
    map ( always [,] map-kind: variable-list ) */
 
@@ -11702,6 +11702,7 @@  c_parser_omp_clause_map (c_parser *parse
 		  || strcmp ("to", p) == 0
 		  || strcmp ("from", p) == 0
 		  || strcmp ("tofrom", p) == 0
+		  || strcmp ("release", p) == 0
 		  || strcmp ("delete", p) == 0)
 		{
 		  c_parser_consume_token (parser);
@@ -11718,13 +11719,15 @@  c_parser_omp_clause_map (c_parser *parse
       if (strcmp ("alloc", p) == 0)
 	kind = GOMP_MAP_ALLOC;
       else if (strcmp ("to", p) == 0)
-	kind = GOMP_MAP_TO;
+	kind = always ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO;
       else if (strcmp ("from", p) == 0)
-	kind = GOMP_MAP_FROM;
+	kind = always ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM;
       else if (strcmp ("tofrom", p) == 0)
-	kind = GOMP_MAP_TOFROM;
+	kind = always ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM;
+      else if (strcmp ("release", p) == 0)
+	kind = GOMP_MAP_RELEASE;
       else if (strcmp ("delete", p) == 0)
-	kind = GOMP_MAP_FORCE_DEALLOC;
+	kind = GOMP_MAP_DELETE;
       else
 	{
 	  c_parser_error (parser, "invalid map kind");
@@ -11732,8 +11735,6 @@  c_parser_omp_clause_map (c_parser *parse
 				     "expected %<)%>");
 	  return list;
 	}
-      if (always)
-	kind = (enum gomp_map_kind) (kind | GOMP_MAP_FLAG_FORCE);
       c_parser_consume_token (parser);
       c_parser_consume_token (parser);
     }
@@ -14237,11 +14238,40 @@  c_parser_omp_target_data (location_t loc
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
 				"#pragma omp target data");
-  if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
+  int map_seen = 0;
+  for (tree *pc = &clauses; *pc;)
     {
-      error_at (loc,
-		"%<#pragma omp target data%> must contain at least one "
-		"%<map%> clause");
+      if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
+	switch (OMP_CLAUSE_MAP_KIND (*pc))
+	  {
+	  case GOMP_MAP_TO:
+	  case GOMP_MAP_ALWAYS_TO:
+	  case GOMP_MAP_FROM:
+	  case GOMP_MAP_ALWAYS_FROM:
+	  case GOMP_MAP_TOFROM:
+	  case GOMP_MAP_ALWAYS_TOFROM:
+	  case GOMP_MAP_ALLOC:
+	  case GOMP_MAP_POINTER:
+	    map_seen = 3;
+	    break;
+	  default:
+	    map_seen |= 1;
+	    error_at (OMP_CLAUSE_LOCATION (*pc),
+		      "%<#pragma omp target data%> with map-type other "
+		      "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
+		      "on %<map%> clause");
+	    *pc = OMP_CLAUSE_CHAIN (*pc);
+	    continue;
+	  }
+      pc = &OMP_CLAUSE_CHAIN (*pc);
+    }
+
+  if (map_seen != 3)
+    {
+      if (map_seen == 0)
+	error_at (loc,
+		  "%<#pragma omp target data%> must contain at least "
+		  "one %<map%> clause");
       return NULL_TREE;
     }
 
@@ -14348,10 +14378,12 @@  c_parser_omp_target_enter_data (location
   for (tree *pc = &clauses; *pc;)
     {
       if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
-	switch (OMP_CLAUSE_MAP_KIND (*pc) & ~GOMP_MAP_FLAG_FORCE)
+	switch (OMP_CLAUSE_MAP_KIND (*pc))
 	  {
 	  case GOMP_MAP_TO:
+	  case GOMP_MAP_ALWAYS_TO:
 	  case GOMP_MAP_ALLOC:
+	  case GOMP_MAP_POINTER:
 	    map_seen = 3;
 	    break;
 	  default:
@@ -14430,17 +14462,21 @@  c_parser_omp_target_exit_data (location_
   for (tree *pc = &clauses; *pc;)
     {
       if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
-	switch (OMP_CLAUSE_MAP_KIND (*pc) & ~GOMP_MAP_FLAG_FORCE)
+	switch (OMP_CLAUSE_MAP_KIND (*pc))
 	  {
 	  case GOMP_MAP_FROM:
-	  case GOMP_MAP_FORCE_DEALLOC & ~GOMP_MAP_FLAG_FORCE:
+	  case GOMP_MAP_ALWAYS_FROM:
+	  case GOMP_MAP_RELEASE:
+	  case GOMP_MAP_DELETE:
+	  case GOMP_MAP_POINTER:
 	    map_seen = 3;
 	    break;
 	  default:
 	    map_seen |= 1;
 	    error_at (OMP_CLAUSE_LOCATION (*pc),
 		      "%<#pragma omp target exit data%> with map-type other "
-		      "than %<from%> or %<delete%> on %<map%> clause");
+		      "than %<from%>, %<release> or %<delete%> on %<map%>"
+		      " clause");
 	    *pc = OMP_CLAUSE_CHAIN (*pc);
 	    continue;
 	  }
@@ -14480,6 +14516,7 @@  c_parser_omp_target (c_parser *parser, e
 {
   location_t loc = c_parser_peek_token (parser)->location;
   c_parser_consume_pragma (parser);
+  tree *pc = NULL, stmt, block;
 
   if (context != pragma_stmt && context != pragma_compound)
     {
@@ -14519,7 +14556,8 @@  c_parser_omp_target (c_parser *parser, e
 	  OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
 	  OMP_TARGET_BODY (stmt) = block;
 	  add_stmt (stmt);
-	  return true;
+	  pc = &OMP_TARGET_CLAUSES (stmt);
+	  goto check_clauses;
 	}
       else if (!flag_openmp)  /* flag_openmp_simd  */
 	{
@@ -14551,19 +14589,46 @@  c_parser_omp_target (c_parser *parser, e
 	}
     }
 
-  tree stmt = make_node (OMP_TARGET);
+  stmt = make_node (OMP_TARGET);
   TREE_TYPE (stmt) = void_type_node;
 
   OMP_TARGET_CLAUSES (stmt)
     = c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
 				"#pragma omp target");
+  pc = &OMP_TARGET_CLAUSES (stmt);
   keep_next_level ();
-  tree block = c_begin_compound_stmt (true);
+  block = c_begin_compound_stmt (true);
   add_stmt (c_parser_omp_structured_block (parser));
   OMP_TARGET_BODY (stmt) = c_end_compound_stmt (loc, block, true);
 
   SET_EXPR_LOCATION (stmt, loc);
   add_stmt (stmt);
+
+check_clauses:
+  while (*pc)
+    {
+      if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
+	switch (OMP_CLAUSE_MAP_KIND (*pc))
+	  {
+	  case GOMP_MAP_TO:
+	  case GOMP_MAP_ALWAYS_TO:
+	  case GOMP_MAP_FROM:
+	  case GOMP_MAP_ALWAYS_FROM:
+	  case GOMP_MAP_TOFROM:
+	  case GOMP_MAP_ALWAYS_TOFROM:
+	  case GOMP_MAP_ALLOC:
+	  case GOMP_MAP_POINTER:
+	    break;
+	  default:
+	    error_at (OMP_CLAUSE_LOCATION (*pc),
+		      "%<#pragma omp target%> with map-type other "
+		      "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
+		      "on %<map%> clause");
+	    *pc = OMP_CLAUSE_CHAIN (*pc);
+	    continue;
+	  }
+      pc = &OMP_CLAUSE_CHAIN (*pc);
+    }
   return true;
 }
 
--- gcc/cp/parser.c.jj	2015-06-03 19:48:09.000000000 +0200
+++ gcc/cp/parser.c	2015-06-11 13:03:19.746878152 +0200
@@ -29150,7 +29150,7 @@  cp_parser_omp_clause_depend (cp_parser *
 
    OpenMP 4.1:
    map-kind:
-     alloc | to | from | tofrom | delete
+     alloc | to | from | tofrom | release | delete
 
    map ( always [,] map-kind: variable-list ) */
 
@@ -29195,13 +29195,15 @@  cp_parser_omp_clause_map (cp_parser *par
       if (strcmp ("alloc", p) == 0)
 	kind = GOMP_MAP_ALLOC;
       else if (strcmp ("to", p) == 0)
-	kind = GOMP_MAP_TO;
+	kind = always ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO;
       else if (strcmp ("from", p) == 0)
-	kind = GOMP_MAP_FROM;
+	kind = always ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM;
       else if (strcmp ("tofrom", p) == 0)
-	kind = GOMP_MAP_TOFROM;
+	kind = always ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM;
+      else if (strcmp ("release", p) == 0)
+	kind = GOMP_MAP_RELEASE;
       else if (strcmp ("delete", p) == 0)
-	kind = GOMP_MAP_FORCE_DEALLOC;
+	kind = GOMP_MAP_DELETE;
       else
 	{
 	  cp_parser_error (parser, "invalid map kind");
@@ -29210,8 +29212,6 @@  cp_parser_omp_clause_map (cp_parser *par
 						 /*consume_paren=*/true);
 	  return list;
 	}
-      if (always)
-	kind = (enum gomp_map_kind) (kind | GOMP_MAP_FLAG_FORCE);
       cp_lexer_consume_token (parser->lexer);
       cp_lexer_consume_token (parser->lexer);
     }
@@ -31690,11 +31690,40 @@  cp_parser_omp_target_data (cp_parser *pa
   tree clauses
     = cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
 				 "#pragma omp target data", pragma_tok);
-  if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
+  int map_seen = 0;
+  for (tree *pc = &clauses; *pc;)
     {
-      error_at (pragma_tok->location,
-		"%<#pragma omp target data%> must contain at least one "
-		"%<map%> clause");
+      if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
+       switch (OMP_CLAUSE_MAP_KIND (*pc))
+	 {
+	 case GOMP_MAP_TO:
+	 case GOMP_MAP_ALWAYS_TO:
+	 case GOMP_MAP_FROM:
+	 case GOMP_MAP_ALWAYS_FROM:
+	 case GOMP_MAP_TOFROM:
+	 case GOMP_MAP_ALWAYS_TOFROM:
+	 case GOMP_MAP_ALLOC:
+	 case GOMP_MAP_POINTER:
+	   map_seen = 3;
+	   break;
+	 default:
+	   map_seen |= 1;
+	   error_at (OMP_CLAUSE_LOCATION (*pc),
+		     "%<#pragma omp target data%> with map-type other "
+		     "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
+		     "on %<map%> clause");
+	   *pc = OMP_CLAUSE_CHAIN (*pc);
+	   continue;
+	 }
+      pc = &OMP_CLAUSE_CHAIN (*pc);
+    }
+
+  if (map_seen != 3)
+    {
+      if (map_seen == 0)
+	error_at (pragma_tok->location,
+		  "%<#pragma omp target data%> must contain at least "
+		  "one %<map%> clause");
       return NULL_TREE;
     }
 
@@ -31759,10 +31788,12 @@  cp_parser_omp_target_enter_data (cp_pars
   for (tree *pc = &clauses; *pc;)
     {
       if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
-       switch (OMP_CLAUSE_MAP_KIND (*pc) & ~GOMP_MAP_FLAG_FORCE)
+       switch (OMP_CLAUSE_MAP_KIND (*pc))
 	 {
 	 case GOMP_MAP_TO:
+	 case GOMP_MAP_ALWAYS_TO:
 	 case GOMP_MAP_ALLOC:
+	 case GOMP_MAP_POINTER:
 	   map_seen = 3;
 	   break;
 	 default:
@@ -31842,17 +31873,21 @@  cp_parser_omp_target_exit_data (cp_parse
   for (tree *pc = &clauses; *pc;)
     {
       if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
-       switch (OMP_CLAUSE_MAP_KIND (*pc) & ~GOMP_MAP_FLAG_FORCE)
+       switch (OMP_CLAUSE_MAP_KIND (*pc))
 	 {
 	 case GOMP_MAP_FROM:
-	 case GOMP_MAP_FORCE_DEALLOC & ~GOMP_MAP_FLAG_FORCE:
+	 case GOMP_MAP_ALWAYS_FROM:
+	 case GOMP_MAP_RELEASE:
+	 case GOMP_MAP_DELETE:
+	 case GOMP_MAP_POINTER:
 	   map_seen = 3;
 	   break;
 	 default:
 	   map_seen |= 1;
 	   error_at (OMP_CLAUSE_LOCATION (*pc),
 		     "%<#pragma omp target exit data%> with map-type other "
-		     "than %<from%> or %<delete%> on %<map%> clause");
+		     "than %<from%>, %<release%> or %<delete%> on %<map%>"
+		     " clause");
 	   *pc = OMP_CLAUSE_CHAIN (*pc);
 	   continue;
 	 }
@@ -31934,6 +31969,8 @@  static bool
 cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
 		      enum pragma_context context)
 {
+  tree *pc = NULL, stmt;
+
   if (context != pragma_stmt && context != pragma_compound)
     {
       cp_parser_error (parser, "expected declaration specifiers");
@@ -31975,7 +32012,8 @@  cp_parser_omp_target (cp_parser *parser,
 	  OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
 	  OMP_TARGET_BODY (stmt) = body;
 	  add_stmt (stmt);
-	  return true;
+	  pc = &OMP_TARGET_CLAUSES (stmt);
+	  goto check_clauses;
 	}
       else if (!flag_openmp)  /* flag_openmp_simd  */
 	{
@@ -32007,17 +32045,44 @@  cp_parser_omp_target (cp_parser *parser,
 	}
     }
 
-  tree stmt = make_node (OMP_TARGET);
+  stmt = make_node (OMP_TARGET);
   TREE_TYPE (stmt) = void_type_node;
 
   OMP_TARGET_CLAUSES (stmt)
     = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
 				 "#pragma omp target", pragma_tok);
+  pc = &OMP_TARGET_CLAUSES (stmt);
   keep_next_level (true);
   OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser);
 
   SET_EXPR_LOCATION (stmt, pragma_tok->location);
   add_stmt (stmt);
+
+check_clauses:
+  while (*pc)
+    {
+      if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
+	switch (OMP_CLAUSE_MAP_KIND (*pc))
+	  {
+	  case GOMP_MAP_TO:
+	  case GOMP_MAP_ALWAYS_TO:
+	  case GOMP_MAP_FROM:
+	  case GOMP_MAP_ALWAYS_FROM:
+	  case GOMP_MAP_TOFROM:
+	  case GOMP_MAP_ALWAYS_TOFROM:
+	  case GOMP_MAP_ALLOC:
+	  case GOMP_MAP_POINTER:
+	    break;
+	  default:
+	    error_at (OMP_CLAUSE_LOCATION (*pc),
+		      "%<#pragma omp target%> with map-type other "
+		      "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
+		      "on %<map%> clause");
+	    *pc = OMP_CLAUSE_CHAIN (*pc);
+	    continue;
+	  }
+      pc = &OMP_CLAUSE_CHAIN (*pc);
+    }
   return true;
 }