https://gcc.gnu.org/g:09e74059bdb12c784570d40b28a33f22090badfe
commit r17-680-g09e74059bdb12c784570d40b28a33f22090badfe Author: Julian Brown <[email protected]> Date: Fri May 22 23:55:50 2026 +0200 OpenMP: Enable 'declare mapper' mappers for 'target update' directives This patch enables use of 'declare mapper' for 'target update' directives, for C and C++ but not yet for Fortran. There are some implementation choices here and some "read-between-the-lines" consequences regarding this functionality, as follows: * It is possible to invoke a mapper which contains clauses that don't make sense for a given 'target update' operation. E.g. if a mapper definition specifies a "from:" mapping and the user does "target update to(...)" which triggers that mapper, the resulting map kind (OpenMP 5.2, "Table 5.3: Map-Type Decay of Map Type Combinations") is "alloc" (and for the inverse case "release"). For such cases, an '-Wopenmp' warning is issued and the map clause in question is dropped from the mapper expansion. (Other choices might be to make this an error, or to do the same thing but silently, or warn only given some special option.) gcc/c-family/ChangeLog: * c-common.h (enum c_omp_region_type): Add C_ORT_UPDATE and C_ORT_OMP_UPDATE codes. * c-omp.cc (omp_basic_map_kind_name): New function. (omp_instantiate_mapper): Add LOC parameter and 'target update' support. (c_omp_instantiate_mappers): Add 'target update' support. gcc/c/ChangeLog: * c-parser.cc (c_parser_omp_clause_from_to): Add parsing for mapper modifier. (c_parser_omp_target_update): Instantiate mappers. gcc/cp/ChangeLog: * parser.cc (cp_parser_omp_clause_from_to): Add parsing for mapper modifier. (cp_parser_omp_target_update): Instantiate mappers. gcc/testsuite/ChangeLog: * c-c++-common/gomp/target-update-iterators-1.c: Update for dg-error change. * c-c++-common/gomp/declare-mapper-19.c: New test. Co-Authored-By: Tobias Burnus <[email protected]> Diff: --- gcc/c-family/c-common.h | 2 + gcc/c-family/c-omp.cc | 126 ++++++++++++++++++--- gcc/c/c-parser.cc | 86 +++++++++++++- gcc/cp/parser.cc | 97 +++++++++++++++- .../c-c++-common/gomp/declare-mapper-19.c | 38 +++++++ .../c-c++-common/gomp/target-update-iterators-1.c | 2 +- 6 files changed, 323 insertions(+), 28 deletions(-) diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h index a4cd67861092..e448868ac1e8 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1312,9 +1312,11 @@ enum c_omp_region_type 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 diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc index b48243a22616..c6bdcf8244c3 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -4562,12 +4562,39 @@ omp_map_decayed_kind (enum gomp_map_kind mapper_kind, 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) { @@ -4600,7 +4627,6 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr, 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) @@ -4626,9 +4652,13 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr, 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); @@ -4646,7 +4676,7 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr, = 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; } @@ -4657,8 +4687,52 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr, 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; @@ -4676,17 +4750,25 @@ c_omp_instantiate_mappers (tree clauses, enum c_omp_region_type ort) 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); @@ -4723,9 +4805,22 @@ c_omp_instantiate_mappers (tree clauses, enum c_omp_region_type ort) 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); @@ -4742,7 +4837,8 @@ c_omp_instantiate_mappers (tree clauses, enum c_omp_region_type ort) { 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) diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index fd6a10d86d0d..7e58d9715988 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -1,6 +1,6 @@ +#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. @@ -21255,7 +21255,7 @@ c_parser_omp_clause_device_type (c_parser *parser, tree list) 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, @@ -21297,8 +21297,9 @@ 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) { @@ -21331,11 +21332,59 @@ c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind, 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; @@ -21365,6 +21414,27 @@ c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind, 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; } @@ -27037,7 +27107,11 @@ c_parser_omp_target_update (location_t loc, c_parser *parser, 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) { diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 0c78c437019f..b5c8c62118cc 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -45179,7 +45179,7 @@ cp_parser_omp_clause_doacross (cp_parser *parser, tree list, location_t loc) 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, @@ -45218,8 +45218,9 @@ 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) { @@ -45259,12 +45260,70 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind, 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, @@ -45295,6 +45354,27 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind, 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; } @@ -51925,7 +52005,12 @@ cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok, 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) { diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-19.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-19.c new file mode 100644 index 000000000000..b7821486fc0c --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-19.c @@ -0,0 +1,38 @@ +/* { 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; +} diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterators-1.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-1.c index d69fe951fed7..37dd69e15d1e 100644 --- a/gcc/testsuite/c-c++-common/gomp/target-update-iterators-1.c +++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-1.c @@ -15,7 +15,7 @@ void f (int **x, float **y) #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 } */ }
