From: Jakub Jelinek <jakub@redhat.com>
To: Ilya Verbin <iverbin@gmail.com>,
Thomas Schwinge <thomas@codesourcery.com>
Cc: gcc-patches@gcc.gnu.org, Kirill Yukhin <kirill.yukhin@gmail.com>
Subject: [gomp4.1] map clause parsing improvements
Date: Thu, 11 Jun 2015 12:52:00 -0000 [thread overview]
Message-ID: <20150611121420.GY10247@tucnak.redhat.com> (raw)
In-Reply-To: <20150609183608.GA47936@msticlxl57.ims.intel.com>
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.
--- 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;
}
Jakub
next prev parent reply other threads:[~2015-06-11 12:14 UTC|newest]
Thread overview: 33+ messages / expand[flat|nested] mbox.gz Atom feed top
2015-04-29 11:44 [gomp4.1] Initial support for some OpenMP 4.1 construct parsing Jakub Jelinek
2015-04-29 11:55 ` Thomas Schwinge
2015-04-29 12:31 ` Jakub Jelinek
2015-04-29 15:20 ` Thomas Schwinge
2015-06-09 18:39 ` Ilya Verbin
2015-06-09 20:25 ` Jakub Jelinek
2015-06-25 19:47 ` Ilya Verbin
2015-06-25 20:31 ` Jakub Jelinek
2015-07-17 16:47 ` Ilya Verbin
2015-07-17 16:54 ` Jakub Jelinek
2015-07-20 16:18 ` Jakub Jelinek
2015-07-20 18:31 ` Jakub Jelinek
2015-07-23 0:50 ` Jakub Jelinek
2015-07-24 20:33 ` Jakub Jelinek
2015-07-29 17:30 ` [gomp4.1] Various accelerator updates from OpenMP 4.1 Jakub Jelinek
2015-09-04 18:17 ` Ilya Verbin
2015-09-04 18:25 ` Jakub Jelinek
2015-09-07 12:48 ` Jakub Jelinek
2015-07-20 19:40 ` [gomp4.1] Initial support for some OpenMP 4.1 construct parsing Ilya Verbin
2015-08-24 12:38 ` Jakub Jelinek
2015-08-24 19:10 ` Ilya Verbin
2015-06-11 12:52 ` Jakub Jelinek [this message]
2015-10-19 10:34 ` [gomp4.1] map clause parsing improvements Thomas Schwinge
2015-10-19 10:46 ` Jakub Jelinek
2015-10-19 15:14 ` Thomas Schwinge
2015-10-20 10:10 ` Jakub Jelinek
2015-10-26 13:04 ` Ilya Verbin
2015-10-26 13:17 ` Jakub Jelinek
2015-10-26 14:16 ` Ilya Verbin
2016-03-17 14:34 ` Thomas Schwinge
2016-03-17 14:37 ` Jakub Jelinek
2016-03-17 14:55 ` Jakub Jelinek
2016-03-17 15:13 ` Rename GOMP_MAP_FORCE_DEALLOC to GOMP_MAP_DELETE (was: [gomp4.1] map clause parsing improvements) Thomas Schwinge
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=20150611121420.GY10247@tucnak.redhat.com \
--to=jakub@redhat.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=iverbin@gmail.com \
--cc=kirill.yukhin@gmail.com \
--cc=thomas@codesourcery.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).