C_ORT_EXIT_DATA = 1 << 4,
C_ORT_INTEROP = 1 << 5,
C_ORT_DECLARE_MAPPER = 1 << 6,
+ C_ORT_UPDATE = 1 << 7,
C_ORT_OMP_DECLARE_SIMD = C_ORT_OMP | C_ORT_DECLARE_SIMD,
C_ORT_OMP_TARGET = C_ORT_OMP | C_ORT_TARGET,
C_ORT_OMP_EXIT_DATA = C_ORT_OMP | C_ORT_EXIT_DATA,
+ C_ORT_OMP_UPDATE = C_ORT_OMP | C_ORT_UPDATE,
C_ORT_OMP_INTEROP = C_ORT_OMP | C_ORT_INTEROP,
C_ORT_OMP_DECLARE_MAPPER = C_ORT_OMP | C_ORT_DECLARE_MAPPER,
C_ORT_ACC_TARGET = C_ORT_ACC | C_ORT_TARGET
m_always_p | always_p, m_present_p | present_p);
}
-/* Instantiate a mapper MAPPER for expression EXPR, adding new clauses to
- OUTLIST. OUTER_KIND is the mapping kind to use if not already specified in
- the mapper declaration. */
+/* Return a name to use for a "basic" map kind, e.g. as output from
+ omp_split_map_kind above. */
+
+static const char *
+omp_basic_map_kind_name (enum gomp_map_kind kind)
+{
+ switch (kind)
+ {
+ case GOMP_MAP_ALLOC:
+ return "alloc";
+ case GOMP_MAP_TO:
+ return "to";
+ case GOMP_MAP_FROM:
+ return "from";
+ case GOMP_MAP_TOFROM:
+ return "tofrom";
+ case GOMP_MAP_RELEASE:
+ return "release";
+ case GOMP_MAP_DELETE:
+ return "delete";
+ default:
+ gcc_unreachable ();
+ }
+}
+
+/* Instantiate a mapper MAPPER for expression EXPR at location LOC,
+ adding new clauses to OUTLIST. OUTER_KIND is the mapping kind
+ to use if not already specified in the mapper declaration; ORT
+ denotes the directive that uses the mapper.
+ The resulting OUTLIST is also returned by the function. */
static tree *
-omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,
+omp_instantiate_mapper (location_t loc, tree *outlist, tree mapper, tree expr,
enum gomp_map_kind outer_kind,
enum c_omp_region_type ort)
{
if (TREE_CODE (t) == OMP_ARRAY_SECTION)
{
- location_t loc = OMP_CLAUSE_LOCATION (c);
tree t2 = lang_hooks.decls.omp_map_array_section (loc, t);
if (t2 == t)
walk_tree (&unshared, remap_mapper_decl_1, &map_info, NULL);
+ OMP_CLAUSE_LOCATION (unshared) = loc;
+
enum gomp_map_kind decayed_kind
= omp_map_decayed_kind (clause_kind, outer_kind,
- (ort & C_ORT_EXIT_DATA) != 0);
+ (ort & C_ORT_EXIT_DATA) != 0
+ || (outer_kind == GOMP_MAP_FROM
+ && (ort & C_ORT_UPDATE) != 0));
OMP_CLAUSE_SET_MAP_KIND (unshared, decayed_kind);
type = TYPE_MAIN_VARIANT (type);
= lang_hooks.decls.omp_extract_mapper_directive (mapper_fn);
if (nested_mapper != mapper)
{
- outlist = omp_instantiate_mapper (outlist, nested_mapper,
+ outlist = omp_instantiate_mapper (loc, outlist, nested_mapper,
t, outer_kind, ort);
continue;
}
continue;
}
- *outlist = unshared;
- outlist = &OMP_CLAUSE_CHAIN (unshared);
+ if (ort & C_ORT_UPDATE)
+ {
+ bool force_p, always_p, present_p;
+ decayed_kind
+ = omp_split_map_kind (decayed_kind, &force_p, &always_p,
+ &present_p);
+ /* We don't expect to see these flags here. */
+ gcc_assert (!force_p && !always_p);
+ /* For a "target update" operation, we want to turn the map node
+ expanded from the mapper back into a OMP_CLAUSE_TO or
+ OMP_CLAUSE_FROM node. If we can do neither, emit a warning and
+ drop the clause. */
+ switch (decayed_kind)
+ {
+ case GOMP_MAP_TO:
+ case GOMP_MAP_FROM:
+ {
+ tree xfer
+ = build_omp_clause (loc, (decayed_kind == GOMP_MAP_TO
+ ? OMP_CLAUSE_TO : OMP_CLAUSE_FROM));
+ OMP_CLAUSE_DECL (xfer) = OMP_CLAUSE_DECL (unshared);
+ OMP_CLAUSE_SIZE (xfer) = OMP_CLAUSE_SIZE (unshared);
+ /* For FROM/TO clauses, "present" is represented by a flag.
+ Set it for the expanded clause here. */
+ if (present_p)
+ OMP_CLAUSE_MOTION_PRESENT (xfer) = 1;
+ *outlist = xfer;
+ outlist = &OMP_CLAUSE_CHAIN (xfer);
+ }
+ break;
+ default:
+ clause_kind
+ = omp_split_map_kind (clause_kind, &force_p, &always_p,
+ &present_p);
+ warning_at (loc, OPT_Wopenmp,
+ "dropping %qs clause during mapper expansion "
+ "in %<#pragma omp target update%>",
+ omp_basic_map_kind_name (clause_kind));
+ inform (OMP_CLAUSE_LOCATION (c), "for map clause here");
+ }
+ }
+ else
+ {
+ *outlist = unshared;
+ outlist = &OMP_CLAUSE_CHAIN (unshared);
+ }
}
return outlist;
for (pc = &clauses, c = clauses; c; c = *pc)
{
bool using_mapper = false;
+ bool update_p = false, update_present_p = false;
switch (OMP_CLAUSE_CODE (c))
{
+ case OMP_CLAUSE_TO:
+ case OMP_CLAUSE_FROM:
+ update_p = true;
+ if (OMP_CLAUSE_MOTION_PRESENT (c))
+ update_present_p = true;
+ /* Fallthrough. */
case OMP_CLAUSE_MAP:
{
tree t = OMP_CLAUSE_DECL (c);
tree type = NULL_TREE;
bool nonunit_array_with_mapper = false;
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME)
+ if (!update_p
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME))
{
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME)
mapper_name = OMP_CLAUSE_DECL (c);
continue;
}
- enum gomp_map_kind kind = OMP_CLAUSE_MAP_KIND (c);
- if (kind == GOMP_MAP_UNSET)
- kind = GOMP_MAP_TOFROM;
+ enum gomp_map_kind kind;
+ if (update_p)
+ {
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO)
+ kind = update_present_p ? GOMP_MAP_PRESENT_TO
+ : GOMP_MAP_TO;
+ else
+ kind = update_present_p ? GOMP_MAP_PRESENT_FROM
+ : GOMP_MAP_FROM;
+ }
+ else
+ {
+ kind = OMP_CLAUSE_MAP_KIND (c);
+ if (kind == GOMP_MAP_UNSET)
+ kind = GOMP_MAP_TOFROM;
+ }
type = TYPE_MAIN_VARIANT (type);
{
tree mapper
= lang_hooks.decls.omp_extract_mapper_directive (mapper_fn);
- pc = omp_instantiate_mapper (pc, mapper, t, kind, ort);
+ pc = omp_instantiate_mapper (OMP_CLAUSE_LOCATION (c),
+ pc, mapper, t, kind, ort);
using_mapper = true;
}
else if (mapper_name)
+#pragma GCC optimize("O0")
/* Parser for C and Objective-C.
Copyright (C) 1987-2026 Free Software Foundation, Inc.
-
Parser actions based on the old Bison parser; structure somewhat
influenced by and fragments based on the C++ parser.
to ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list )
motion-modifier:
- present | iterator (iterators-definition) */
+ present | iterator (iterators-definition) | mapper (id) */
static tree
c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind,
}
}
- bool present = false;
+ bool present = false, mapper_modifier = false;
tree iterators = NULL_TREE;
+ tree mapper_name = NULL_TREE;
for (int pos = 1; pos < colon_pos; ++pos)
{
iterators = c_parser_omp_iterators (parser);
pos += iterator_length - 1;
}
+ else if (strcmp ("mapper", p) == 0)
+ {
+ c_parser_consume_token (parser);
+
+ matching_parens mparens;
+ if (mparens.require_open (parser))
+ {
+ if (mapper_modifier)
+ {
+ c_parser_error (parser, "too many %<mapper%> modifiers");
+ /* Assume it's a well-formed mapper modifier, even if it
+ seems to be in the wrong place. */
+ c_parser_consume_token (parser);
+ mparens.require_close (parser);
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+
+ token = c_parser_peek_token (parser);
+ switch (token->type)
+ {
+ case CPP_NAME:
+ {
+ mapper_name = token->value;
+ c_parser_consume_token (parser);
+ }
+ break;
+ case CPP_KEYWORD:
+ if (token->keyword == RID_DEFAULT)
+ {
+ c_parser_consume_token (parser);
+ break;
+ }
+ /* FALLTHROUGH */
+ default:
+ error_at (token->location,
+ "expected identifier or %<default%>");
+ return list;
+ }
+ if (!mparens.require_close (parser))
+ {
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+ mapper_modifier = true;
+ pos += 3;
+ }
+ }
else
{
error_at (token->location,
- "%qs clause with modifier other than %<iterator%> or "
- "%<present%>",
+ "%qs clause with modifier other than %<iterator%>, "
+ "%<mapper%>, or %<present%>",
kind == OMP_CLAUSE_TO ? "to" : "from");
parens.skip_until_found_close (parser);
return list;
for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
OMP_CLAUSE_ITERATORS (c) = iterators;
+ if (mapper_name)
+ {
+ tree last_new = NULL_TREE;
+ for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+ last_new = c;
+
+ tree name = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_PUSH_MAPPER_NAME);
+ OMP_CLAUSE_DECL (name) = mapper_name;
+ OMP_CLAUSE_CHAIN (name) = nl;
+ nl = name;
+
+ gcc_assert (last_new);
+
+ name = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_POP_MAPPER_NAME);
+ OMP_CLAUSE_DECL (name) = null_pointer_node;
+ OMP_CLAUSE_CHAIN (name) = OMP_CLAUSE_CHAIN (last_new);
+ OMP_CLAUSE_CHAIN (last_new) = name;
+ }
+
return nl;
}
tree clauses
= c_parser_omp_all_clauses (parser, OMP_TARGET_UPDATE_CLAUSE_MASK,
- "#pragma omp target update");
+ "#pragma omp target update",
+ /* finish_p = */ false);
+ clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_UPDATE);
+ clauses = c_finish_omp_clauses (clauses, C_ORT_OMP_UPDATE);
+
if (omp_find_clause (clauses, OMP_CLAUSE_TO) == NULL_TREE
&& omp_find_clause (clauses, OMP_CLAUSE_FROM) == NULL_TREE)
{
to ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list )
motion-modifier:
- present | iterator (iterators-definition) */
+ present | iterator (iterators-definition) | mapper (id) */
static tree
cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind,
}
}
- bool present = false;
+ bool present = false, mapper_modifier = false;
tree iterators = NULL_TREE;
+ tree mapper_name = NULL_TREE;
for (int pos = 1; pos < colon_pos; ++pos)
{
iterators = cp_parser_omp_iterators (parser);
pos += iterator_length - 1;
}
-
+ else if (strcmp ("mapper", p) == 0)
+ {
+ cp_lexer_consume_token (parser->lexer);
+ matching_parens parens;
+ if (parens.require_open (parser))
+ {
+ if (mapper_modifier)
+ {
+ cp_parser_error (parser, "too many %<mapper%> modifiers");
+ /* Assume it's a well-formed mapper modifier, even if it
+ seems to be in the wrong place. */
+ cp_lexer_consume_token (parser->lexer);
+ parens.require_close (parser);
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/
+ true);
+ return list;
+ }
+ token = cp_lexer_peek_token (parser->lexer);
+ switch (token->type)
+ {
+ case CPP_NAME:
+ {
+ cp_expr e = cp_parser_identifier (parser);
+ if (e != error_mark_node)
+ mapper_name = e;
+ else
+ goto err;
+ }
+ break;
+ case CPP_KEYWORD:
+ if (token->keyword == RID_DEFAULT)
+ {
+ cp_lexer_consume_token (parser->lexer);
+ break;
+ }
+ /* Fallthrough. */
+ /* FALLTHROUGH. */
+ default:
+ err:
+ cp_parser_error (parser,
+ "expected identifier or %<default%>");
+ return list;
+ }
+ if (!parens.require_close (parser))
+ {
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/
+ true);
+ return list;
+ }
+ mapper_modifier = true;
+ pos += 3;
+ }
+ }
else
{
error_at (token->location,
- "%qs clause with modifier other than %<iterator%> "
- "or %<present%>",
+ "%qs clause with modifier other than %<iterator%>, "
+ "
%<mapper%>, or %<present%>",
kind == OMP_CLAUSE_TO ? "to" : "from");
cp_parser_skip_to_closing_parenthesis (parser,
/*recovering=*/true,
for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
OMP_CLAUSE_ITERATORS (c) = iterators;
+ if (mapper_name)
+ {
+ tree last_new = NULL_TREE;
+ for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+ last_new = c;
+
+ tree name = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_PUSH_MAPPER_NAME);
+ OMP_CLAUSE_DECL (name) = mapper_name;
+ OMP_CLAUSE_CHAIN (name) = nl;
+ nl = name;
+
+ gcc_assert (last_new);
+
+ name = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_POP_MAPPER_NAME);
+ OMP_CLAUSE_DECL (name) = null_pointer_node;
+ OMP_CLAUSE_CHAIN (name) = OMP_CLAUSE_CHAIN (last_new);
+ OMP_CLAUSE_CHAIN (last_new) = name;
+ }
+
return nl;
}
tree clauses
= cp_parser_omp_all_clauses (parser, OMP_TARGET_UPDATE_CLAUSE_MASK,
- "#pragma omp target update", pragma_tok);
+ "#pragma omp target update", pragma_tok,
+ /* finish_p = */ false);
+ if (!processing_template_decl)
+ clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_UPDATE);
+ clauses = finish_omp_clauses (clauses, C_ORT_OMP_UPDATE);
+
if (omp_find_clause (clauses, OMP_CLAUSE_TO) == NULL_TREE
&& omp_find_clause (clauses, OMP_CLAUSE_FROM) == NULL_TREE)
{
--- /dev/null
+/* { dg-do compile { target { c || c++11 } } } */
+/* { dg-additional-options "-fdump-tree-original" } */
+
+typedef struct {
+ int a, b, c, d;
+} S;
+
+#pragma omp declare mapper (S s) map(alloc: s.a) map(to: s.b) map(from: s.c) \
+ map(tofrom: s.d)
+#pragma omp declare mapper (update: S s) map(s.a, s.b, s.c, s.d)
+
+int main()
+{
+ S v;
+#pragma omp target update to(v)
+/* { dg-warning {dropping .from. clause during mapper expansion in .#pragma omp target update.} "" { target *-*-* } .-1 } */
+/* { dg-warning {dropping .alloc. clause during mapper expansion in .#pragma omp target update.} "" { target *-*-* } .-2 } */
+/* { dg-final { scan-tree-dump-times {(?n)update to\(v\.d\) to\(v\.b\)$} 1 "original" } } */
+#pragma omp target update from(v)
+/* { dg-warning {dropping .to. clause during mapper expansion in .#pragma omp target update.} "" { target *-*-* } .-1 } */
+/* { dg-warning {dropping .alloc. clause during mapper expansion in .#pragma omp target update.} "" { target *-*-* } .-2 } */
+/* { dg-final { scan-tree-dump-times {(?n)update from\(v\.d\) from\(v\.c\)$} 1 "original" } } */
+
+#pragma omp target update to(mapper(update): v)
+/* { dg-final { scan-tree-dump-times {(?n)update to\(v\.d\) to\(v\.c\) to\(v\.b\) to\(v\.a\)$} 1 "original" } } */
+#pragma omp target update from(mapper(update): v)
+/* { dg-final { scan-tree-dump-times {(?n)update from\(v\.d\) from\(v\.c\) from\(v\.b\) from\(v\.a\)$} 1 "original" } } */
+
+#pragma omp target update to(present, mapper(update): v)
+/* { dg-final { scan-tree-dump-times {(?n)update to\(present:v\.d\) to\(present:v\.c\) to\(present:v\.b\) to\(present:v\.a\)$} 2 "original" } } */
+#pragma omp target update from(present, mapper(update): v)
+/* { dg-final { scan-tree-dump-times {(?n)update from\(present:v\.d\) from\(present:v\.c\) from\(present:v\.b\) from\(present:v\.a\)$} 2 "original" } } */
+
+#pragma omp target update to(present: v.a, v.b, v.c, v.d)
+#pragma omp target update from(present: v.a, v.b, v.c, v.d)
+
+ return 0;
+}
#pragma omp target update to (iterator(i=0:DIM1), iterator(j=0:DIM2): x[i][j]) /* { dg-error "too many 'iterator' modifiers" } */
/* { dg-error "'#pragma omp target update' must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-1 } */
- #pragma omp target update from (iterator(i=0:DIM1), something: x[i][j]) /* { dg-error "'from' clause with modifier other than 'iterator' or 'present'" } */
+ #pragma omp target update from (iterator(i=0:DIM1), something: x[i][j]) /* { dg-error "'from' clause with modifier other than 'iterator', 'mapper', or 'present'" } */
/* { dg-error "expected '\\)' before 'something'" "" { target c } .-1 } */
/* { dg-error "'#pragma omp target update' must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 } */
}
projects / thirdparty / gcc.git / commitdiff
