@@ -556,6 +556,8 @@ const struct attribute_spec c_common_gnu_attributes[] =
handle_omp_declare_variant_attribute, NULL },
{ "omp declare variant variant", 0, -1, true, false, false, false,
handle_omp_declare_variant_attribute, NULL },
+ { "omp declare variant adjust_args need_device_ptr", 0, -1, true, false, false, false,
+ handle_omp_declare_variant_attribute, NULL },
{ "simd", 0, 1, true, false, false, false,
handle_simd_attribute, NULL },
{ "omp declare target", 0, -1, true, false, false, false,
@@ -4196,8 +4196,8 @@ const struct c_omp_directive c_omp_directives[] = {
C_OMP_DIR_DECLARATIVE, false },
{ "depobj", nullptr, nullptr, PRAGMA_OMP_DEPOBJ,
C_OMP_DIR_STANDALONE, false },
- /* { "dispatch", nullptr, nullptr, PRAGMA_OMP_DISPATCH,
- C_OMP_DIR_CONSTRUCT, false }, */
+ { "dispatch", nullptr, nullptr, PRAGMA_OMP_DISPATCH,
+ C_OMP_DIR_DECLARATIVE, false },
{ "distribute", nullptr, nullptr, PRAGMA_OMP_DISTRIBUTE,
C_OMP_DIR_CONSTRUCT, true },
{ "end", "assumes", nullptr, PRAGMA_OMP_END,
@@ -1526,6 +1526,7 @@ static const struct omp_pragma_def omp_pragmas[] = {
{ "cancellation", PRAGMA_OMP_CANCELLATION_POINT },
{ "critical", PRAGMA_OMP_CRITICAL },
{ "depobj", PRAGMA_OMP_DEPOBJ },
+ { "dispatch", PRAGMA_OMP_DISPATCH },
{ "error", PRAGMA_OMP_ERROR },
{ "end", PRAGMA_OMP_END },
{ "flush", PRAGMA_OMP_FLUSH },
@@ -55,6 +55,7 @@ enum pragma_kind {
PRAGMA_OMP_CRITICAL,
PRAGMA_OMP_DECLARE,
PRAGMA_OMP_DEPOBJ,
+ PRAGMA_OMP_DISPATCH,
PRAGMA_OMP_DISTRIBUTE,
PRAGMA_OMP_ERROR,
PRAGMA_OMP_END,
@@ -132,9 +133,11 @@ enum pragma_omp_clause {
PRAGMA_OMP_CLAUSE_LINK,
PRAGMA_OMP_CLAUSE_MAP,
PRAGMA_OMP_CLAUSE_MERGEABLE,
+ PRAGMA_OMP_CLAUSE_NOCONTEXT,
PRAGMA_OMP_CLAUSE_NOGROUP,
PRAGMA_OMP_CLAUSE_NONTEMPORAL,
PRAGMA_OMP_CLAUSE_NOTINBRANCH,
+ PRAGMA_OMP_CLAUSE_NOVARIANTS,
PRAGMA_OMP_CLAUSE_NOWAIT,
PRAGMA_OMP_CLAUSE_NUM_TASKS,
PRAGMA_OMP_CLAUSE_NUM_TEAMS,
@@ -1733,6 +1733,8 @@ static void c_parser_omp_assumption_clauses (c_parser *, bool);
static void c_parser_omp_allocate (c_parser *);
static void c_parser_omp_assumes (c_parser *);
static bool c_parser_omp_ordered (c_parser *, enum pragma_context, bool *);
+static tree
+c_parser_omp_dispatch (location_t, c_parser *);
static void c_parser_oacc_routine (c_parser *, enum pragma_context);
/* These Objective-C parser functions are only ever called when
@@ -14972,6 +14974,8 @@ c_parser_omp_clause_name (c_parser *parser)
case 'n':
if (!strcmp ("no_create", p))
result = PRAGMA_OACC_CLAUSE_NO_CREATE;
+ else if (!strcmp ("nocontext", p))
+ result = PRAGMA_OMP_CLAUSE_NOCONTEXT;
else if (!strcmp ("nogroup", p))
result = PRAGMA_OMP_CLAUSE_NOGROUP;
else if (!strcmp ("nohost", p))
@@ -14980,6 +14984,8 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_NONTEMPORAL;
else if (!strcmp ("notinbranch", p))
result = PRAGMA_OMP_CLAUSE_NOTINBRANCH;
+ else if (!strcmp ("novariants", p))
+ result = PRAGMA_OMP_CLAUSE_NOVARIANTS;
else if (!strcmp ("nowait", p))
result = PRAGMA_OMP_CLAUSE_NOWAIT;
else if (!strcmp ("num_gangs", p))
@@ -19190,6 +19196,60 @@ c_parser_omp_clause_uniform (c_parser *parser, tree list)
return list;
}
+/* OpenMP 5.1
+ novariants ( scalar-expression ) */
+
+static tree
+c_parser_omp_clause_novariants (c_parser *parser, tree list)
+{
+ matching_parens parens;
+ if (!parens.require_open (parser))
+ return list;
+
+ location_t loc = c_parser_peek_token (parser)->location;
+ c_expr expr = c_parser_expr_no_commas (parser, NULL);
+ tree t = convert_lvalue_to_rvalue (loc, expr, true, true).value;
+ t = c_objc_common_truthvalue_conversion (loc, t);
+ t = c_fully_fold (t, false, NULL);
+ parens.skip_until_found_close (parser);
+
+ check_no_duplicate_clause (list, OMP_CLAUSE_NOVARIANTS, "novariants");
+
+ tree c = build_omp_clause (loc, OMP_CLAUSE_NOVARIANTS);
+ OMP_CLAUSE_NOVARIANTS_EXPR (c) = t;
+ OMP_CLAUSE_CHAIN (c) = list;
+ list = c;
+
+ return list;
+}
+
+/* OpenMP 5.1
+ nocontext ( scalar-expression ) */
+
+static tree
+c_parser_omp_clause_nocontext (c_parser *parser, tree list)
+{
+ matching_parens parens;
+ if (!parens.require_open (parser))
+ return list;
+
+ location_t loc = c_parser_peek_token (parser)->location;
+ c_expr expr = c_parser_expr_no_commas (parser, NULL);
+ tree t = convert_lvalue_to_rvalue (loc, expr, true, true).value;
+ t = c_objc_common_truthvalue_conversion (loc, t);
+ t = c_fully_fold (t, false, NULL);
+ parens.skip_until_found_close (parser);
+
+ check_no_duplicate_clause (list, OMP_CLAUSE_NOCONTEXT, "nocontext");
+
+ tree c = build_omp_clause (loc, OMP_CLAUSE_NOCONTEXT);
+ OMP_CLAUSE_NOCONTEXT_EXPR (c) = t;
+ OMP_CLAUSE_CHAIN (c) = list;
+ list = c;
+
+ return list;
+}
+
/* OpenMP 5.0:
detach ( event-handle ) */
@@ -19801,6 +19861,14 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses);
c_name = "enter";
break;
+ case PRAGMA_OMP_CLAUSE_NOVARIANTS:
+ c_name = "novariants";
+ clauses = c_parser_omp_clause_novariants (parser, clauses);
+ break;
+ case PRAGMA_OMP_CLAUSE_NOCONTEXT:
+ c_name = "nocontext";
+ clauses = c_parser_omp_clause_nocontext (parser, clauses);
+ break;
default:
c_parser_error (parser, "expected an OpenMP clause");
goto saw_error;
@@ -23458,6 +23526,168 @@ c_parser_omp_scope (location_t loc, c_parser *parser, bool *if_p)
return add_stmt (stmt);
}
+// Adapted from c_parser_expr_no_commas
+static tree
+c_parser_omp_dispatch_body (c_parser *parser)
+{
+ struct c_expr lhs, rhs, ret;
+ struct c_expr orig_expr;
+ location_t expr_loc = c_parser_peek_token (parser)->location;
+ source_range tok_range = c_parser_peek_token (parser)->get_range ();
+ location_t sizeof_arg_loc[3];
+ tree sizeof_arg[3];
+ vec<tree, va_gc> *exprlist;
+ vec<location_t> arg_loc = vNULL;
+ vec<tree, va_gc> *origtypes = NULL;
+ unsigned int literal_zero_mask;
+ location_t start;
+ location_t finish;
+
+ lhs = c_parser_conditional_expression (parser, NULL, NULL);
+ if (TREE_CODE (lhs.value) == CALL_EXPR)
+ return lhs.value;
+ else
+ {
+ location_t op_location = c_parser_peek_token (parser)->location;
+ if (!c_parser_require (parser, CPP_EQ, "expected %<=%>"))
+ return error_mark_node;
+
+ /* Parse function name*/
+ if (!c_parser_next_token_is (parser, CPP_NAME))
+ {
+ c_parser_error (parser, "expected a function name");
+ rhs.set_error ();
+ return rhs.value;
+ }
+ expr_loc = c_parser_peek_token (parser)->location;
+ tree id = c_parser_peek_token (parser)->value;
+ c_parser_consume_token (parser);
+ if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
+ return error_mark_node;
+
+ rhs.value = build_external_ref (expr_loc, id, true, &rhs.original_type);
+ set_c_expr_source_range (&rhs, tok_range);
+ /* Parse argument list */
+ {
+ for (int i = 0; i < 3; i++)
+ {
+ sizeof_arg[i] = NULL_TREE;
+ sizeof_arg_loc[i] = UNKNOWN_LOCATION;
+ }
+ literal_zero_mask = 0;
+ if (c_parser_next_token_is (parser, CPP_CLOSE_PAREN))
+ exprlist = NULL;
+ else
+ exprlist = c_parser_expr_list (parser, true, false, &origtypes,
+ sizeof_arg_loc, sizeof_arg, &arg_loc,
+ &literal_zero_mask);
+ c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>");
+ }
+ orig_expr = rhs;
+ mark_exp_read (rhs.value);
+ if (warn_sizeof_pointer_memaccess)
+ sizeof_pointer_memaccess_warning (sizeof_arg_loc, rhs.value, exprlist,
+ sizeof_arg,
+ sizeof_ptr_memacc_comptypes);
+ if (TREE_CODE (rhs.value) == FUNCTION_DECL)
+ {
+ if (fndecl_built_in_p (rhs.value, BUILT_IN_MEMSET)
+ && vec_safe_length (exprlist) == 3)
+ {
+ tree arg0 = (*exprlist)[0];
+ tree arg2 = (*exprlist)[2];
+ warn_for_memset (expr_loc, arg0, arg2, literal_zero_mask);
+ }
+ if (warn_absolute_value
+ && fndecl_built_in_p (rhs.value, BUILT_IN_NORMAL)
+ && vec_safe_length (exprlist) == 1)
+ warn_for_abs (expr_loc, rhs.value, (*exprlist)[0]);
+ if (parser->omp_for_parse_state
+ && parser->omp_for_parse_state->in_intervening_code
+ && omp_runtime_api_call (rhs.value))
+ {
+ error_at (expr_loc, "calls to the OpenMP runtime API are "
+ "not permitted in intervening code");
+ parser->omp_for_parse_state->fail = true;
+ }
+ }
+
+ start = rhs.get_start ();
+ finish = parser->tokens_buf[0].get_finish ();
+ rhs.value = c_build_function_call_vec (expr_loc, arg_loc, rhs.value,
+ exprlist, origtypes);
+ set_c_expr_source_range (&rhs, start, finish);
+ rhs.m_decimal = 0;
+
+ rhs.original_code = ERROR_MARK;
+ if (TREE_CODE (rhs.value) == INTEGER_CST
+ && TREE_CODE (orig_expr.value) == FUNCTION_DECL
+ && fndecl_built_in_p (orig_expr.value, BUILT_IN_CONSTANT_P))
+ rhs.original_code = C_MAYBE_CONST_EXPR;
+ rhs.original_type = NULL;
+ if (exprlist)
+ {
+ release_tree_vector (exprlist);
+ release_tree_vector (origtypes);
+ }
+ arg_loc.release ();
+
+ /* Build assignment */
+ rhs = convert_lvalue_to_rvalue (expr_loc, rhs, true, true);
+ ret.value
+ = build_modify_expr (op_location, lhs.value, lhs.original_type,
+ NOP_EXPR, expr_loc, rhs.value, rhs.original_type);
+ ret.m_decimal = 0;
+ set_c_expr_source_range (&ret, lhs.get_start (), rhs.get_finish ());
+ ret.original_code = MODIFY_EXPR;
+ ret.original_type = NULL;
+ return ret.value;
+ }
+}
+
+/* OpenMP 5.1:
+ # pragma omp dispatch dispatch-clause[optseq] new-line
+ expression-stmt
+
+ LOC is the location of the #pragma.
+*/
+
+#define OMP_DISPATCH_CLAUSE_MASK \
+ ((OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOVARIANTS) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOCONTEXT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT))
+
+static tree
+c_parser_omp_dispatch (location_t loc, c_parser *parser)
+{
+ tree stmt = make_node (OMP_DISPATCH);
+ SET_EXPR_LOCATION (stmt, loc);
+ TREE_TYPE (stmt) = void_type_node;
+
+ OMP_DISPATCH_CLAUSES (stmt)
+ = c_parser_omp_all_clauses (parser, OMP_DISPATCH_CLAUSE_MASK,
+ "#pragma omp dispatch");
+
+ // Parse body as expression statement
+ loc = c_parser_peek_token (parser)->location;
+ tree dispatch_body = c_parser_omp_dispatch_body (parser);
+ if (dispatch_body == error_mark_node)
+ {
+ inform (loc, "%<#pragma omp dispatch%> must be followed by a function "
+ "call with optional assignment");
+ c_parser_skip_to_end_of_block_or_statement (parser);
+ return NULL_TREE;
+ }
+
+ c_parser_skip_until_found (parser, CPP_SEMICOLON, "expected %<;%>");
+ OMP_DISPATCH_BODY (stmt) = dispatch_body;
+
+ return add_stmt (stmt);
+}
+
/* OpenMP 3.0:
# pragma omp task task-clause[optseq] new-line
@@ -24438,6 +24668,10 @@ check_clauses:
OpenMP 5.0:
# pragma omp declare variant (identifier) match(context-selector) new-line
+
+ OpenMP 5.1
+ # pragma omp declare variant (identifier) match(context-selector) \
+ adjust_args(adjust-op:argument-list) new-line
*/
#define OMP_DECLARE_SIMD_CLAUSE_MASK \
@@ -24901,19 +25135,39 @@ c_finish_omp_declare_variant (c_parser *parser, tree fndecl, tree parms)
parens.require_close (parser);
- if (c_parser_next_token_is (parser, CPP_COMMA)
- && c_parser_peek_2nd_token (parser)->type == CPP_NAME)
- c_parser_consume_token (parser);
+ vec<tree> adjust_args_list = vNULL;
+ bool has_match = false, has_adjust_args = false;
+ location_t adjust_args_loc;
- const char *clause = "";
- location_t match_loc = c_parser_peek_token (parser)->location;
- if (c_parser_next_token_is (parser, CPP_NAME))
- clause = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
- if (strcmp (clause, "match"))
+ do
{
- c_parser_error (parser, "expected %<match%>");
- goto fail;
- }
+ if (c_parser_next_token_is (parser, CPP_COMMA)
+ && c_parser_peek_2nd_token (parser)->type == CPP_NAME)
+ c_parser_consume_token (parser);
+
+ const char *clause = "";
+ location_t match_loc = c_parser_peek_token (parser)->location;
+ if (c_parser_next_token_is (parser, CPP_NAME))
+ clause = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+
+ enum clause
+ {
+ match,
+ adjust_args
+ } ccode;
+
+ if (strcmp (clause, "match") == 0)
+ ccode = match;
+ else if (strcmp (clause, "adjust_args") == 0)
+ {
+ ccode = adjust_args;
+ adjust_args_loc = match_loc;
+ }
+ else
+ {
+ c_parser_error (parser, "expected %<match%> or %<adjust_args%>");
+ goto fail;
+ }
c_parser_consume_token (parser);
@@ -24923,55 +25177,172 @@ c_finish_omp_declare_variant (c_parser *parser, tree fndecl, tree parms)
if (parms == NULL_TREE)
parms = error_mark_node;
- tree ctx = c_parser_omp_context_selector_specification (parser, parms);
- if (ctx == error_mark_node)
- goto fail;
- ctx = omp_check_context_selector (match_loc, ctx);
- if (ctx != error_mark_node && variant != error_mark_node)
+ if (ccode == match)
{
- if (TREE_CODE (variant) != FUNCTION_DECL)
+ has_match = true;
+ tree ctx = c_parser_omp_context_selector_specification (parser, parms);
+ if (ctx == error_mark_node)
+ goto fail;
+ ctx = omp_check_context_selector (match_loc, ctx);
+ if (ctx != error_mark_node && variant != error_mark_node)
{
- error_at (token->location, "variant %qD is not a function", variant);
- variant = error_mark_node;
- }
- else if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT,
- OMP_TRAIT_CONSTRUCT_SIMD)
- && !comptypes (TREE_TYPE (fndecl), TREE_TYPE (variant)))
- {
- error_at (token->location, "variant %qD and base %qD have "
- "incompatible types", variant, fndecl);
- variant = error_mark_node;
- }
- else if (fndecl_built_in_p (variant)
- && (strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)),
- "__builtin_", strlen ("__builtin_")) == 0
- || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)),
- "__sync_", strlen ("__sync_")) == 0
- || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)),
- "__atomic_", strlen ("__atomic_")) == 0))
- {
- error_at (token->location, "variant %qD is a built-in", variant);
- variant = error_mark_node;
- }
- if (variant != error_mark_node)
- {
- C_DECL_USED (variant) = 1;
- tree construct
- = omp_get_context_selector_list (ctx, OMP_TRAIT_SET_CONSTRUCT);
- omp_mark_declare_variant (match_loc, variant, construct);
- if (omp_context_selector_matches (ctx))
+ if (TREE_CODE (variant) != FUNCTION_DECL)
{
- tree attr
- = tree_cons (get_identifier ("omp declare variant base"),
- build_tree_list (variant, ctx),
- DECL_ATTRIBUTES (fndecl));
- DECL_ATTRIBUTES (fndecl) = attr;
+ error_at (token->location, "variant %qD is not a function",
+ variant);
+ variant = error_mark_node;
+ }
+ else if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT,
+ OMP_TRAIT_CONSTRUCT_SIMD)
+ && !comptypes (TREE_TYPE (fndecl), TREE_TYPE (variant)))
+ {
+ error_at (token->location,
+ "variant %qD and base %qD have "
+ "incompatible types",
+ variant, fndecl);
+ variant = error_mark_node;
+ }
+ else if (fndecl_built_in_p (variant)
+ && (strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)),
+ "__builtin_", strlen ("__builtin_"))
+ == 0
+ || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)),
+ "__sync_", strlen ("__sync_"))
+ == 0
+ || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)),
+ "__atomic_", strlen ("__atomic_"))
+ == 0))
+ {
+ error_at (token->location, "variant %qD is a built-in", variant);
+ variant = error_mark_node;
+ }
+ if (variant != error_mark_node)
+ {
+ C_DECL_USED (variant) = 1;
+ tree construct
+ = omp_get_context_selector_list (ctx, OMP_TRAIT_SET_CONSTRUCT);
+ omp_mark_declare_variant (match_loc, variant, construct);
+ if (omp_context_selector_matches (ctx))
+ {
+ tree attr
+ = tree_cons (get_identifier ("omp declare variant base"),
+ build_tree_list (variant, ctx),
+ DECL_ATTRIBUTES (fndecl));
+ DECL_ATTRIBUTES (fndecl) = attr;
+ }
}
}
}
+ else if (ccode == adjust_args)
+ {
+ has_adjust_args = true;
+ if (c_parser_next_token_is (parser, CPP_NAME)
+ && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+ {
+ const char *p
+ = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+ if (strcmp (p, "need_device_ptr") == 0
+ || strcmp (p, "nothing") == 0)
+ {
+ c_parser_consume_token (parser); // need_device_ptr
+ c_parser_consume_token (parser); // :
+
+ location_t loc = c_parser_peek_token (parser)->location;
+ tree list
+ = c_parser_omp_variable_list (parser, loc, OMP_CLAUSE_ERROR,
+ NULL_TREE);
+
+ tree arg;
+ for (tree c = list; c != NULL_TREE; c = TREE_CHAIN (c))
+ {
+ tree decl = TREE_PURPOSE (c);
+ int idx;
+ for (arg = parms, idx = 0; arg != NULL;
+ arg = TREE_CHAIN (arg), idx++)
+ if (arg == decl)
+ break;
+ if (arg == NULL_TREE)
+ {
+ error_at (loc, "%qD is not a function argument",
+ decl);
+ goto fail;
+ }
+ if (adjust_args_list.contains (arg))
+ {
+ error_at (loc, "%qD is specified more than once",
+ decl);
+ goto fail;
+ }
+ if (strcmp (p, "need_device_ptr") == 0
+ && TREE_CODE (TREE_TYPE (arg)) != POINTER_TYPE)
+ {
+ error_at (loc, "%qD is not a C pointer", decl);
+ goto fail;
+ }
+ adjust_args_list.safe_push (arg);
+ if (strcmp (p, "need_device_ptr") == 0)
+ {
+ tree variant_decl = (TREE_CODE (variant) == NOP_EXPR)
+ ? TREE_OPERAND (variant, 0)
+ : variant;
+ tree variant_parm = DECL_ARGUMENTS (variant_decl);
+ for (int i = 0; i < idx; i++)
+ {
+ variant_parm = TREE_CHAIN (variant_parm);
+ gcc_assert (variant_parm != NULL);
+ }
+ tree attr = tree_cons (
+ get_identifier ("omp declare variant adjust_args "
+ "need_device_ptr"),
+ NULL_TREE, DECL_ATTRIBUTES (variant_parm));
+ DECL_ATTRIBUTES (variant_parm) = attr;
+ }
+ }
+ }
+ else
+ {
+ error_at (c_parser_peek_token (parser)->location,
+ "expected %<nothing%> or %<need_device_ptr%>");
+ goto fail;
+ }
+ }
+ else
+ {
+ error_at (c_parser_peek_token (parser)->location,
+ "expected %<nothing%> or %<need_device_ptr%> "
+ "followed by %<:%>");
+ goto fail;
+ }
+ }
parens.require_close (parser);
+ } while (c_parser_next_token_is_not (parser, CPP_PRAGMA_EOL));
c_parser_skip_to_pragma_eol (parser);
+
+ if (has_adjust_args)
+ {
+ if (!has_match)
+ {
+ error_at (
+ adjust_args_loc,
+ "an %<adjust_args%> clause can only be specified if the "
+ "%<dispatch%> selector of the construct selector set appears "
+ "in the %<match%> clause");
+ }
+ else
+ {
+ tree attr = lookup_attribute ("omp declare variant base",
+ DECL_ATTRIBUTES (fndecl));
+ tree ctx = TREE_VALUE (TREE_VALUE (attr));
+ if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT,
+ OMP_TRAIT_CONSTRUCT_DISPATCH))
+ error_at (
+ adjust_args_loc,
+ "an %<adjust_args%> clause can only be specified if the "
+ "%<dispatch%> selector of the construct selector set appears "
+ "in the %<match%> clause");
+ }
+ }
}
/* Finalize #pragma omp declare simd or #pragma omp declare variant
@@ -25788,7 +26159,6 @@ c_parser_omp_declare_reduction (c_parser *parser, enum pragma_context context)
types.release ();
}
-
/* OpenMP 4.0
#pragma omp declare simd declare-simd-clauses[optseq] new-line
#pragma omp declare reduction (reduction-id : typename-list : expression) \
@@ -25796,7 +26166,11 @@ c_parser_omp_declare_reduction (c_parser *parser, enum pragma_context context)
#pragma omp declare target new-line
OpenMP 5.0
- #pragma omp declare variant (identifier) match (context-selector) */
+ #pragma omp declare variant (identifier) match (context-selector)
+
+ OpenMP 5.1
+ #pragma omp declare variant (identifier) match (context-selector) \
+ adjust_args(adjust-op:argument-list) */
static bool
c_parser_omp_declare (c_parser *parser, enum pragma_context context)
@@ -26599,6 +26973,9 @@ c_parser_omp_construct (c_parser *parser, bool *if_p)
case PRAGMA_OMP_ASSUME:
c_parser_omp_assume (parser, if_p);
return;
+ case PRAGMA_OMP_DISPATCH:
+ stmt = c_parser_omp_dispatch (loc, parser);
+ break;
default:
gcc_unreachable ();
}
@@ -16054,6 +16054,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
case OMP_CLAUSE_FINALIZE:
case OMP_CLAUSE_NOHOST:
case OMP_CLAUSE_INDIRECT:
+ case OMP_CLAUSE_NOVARIANTS:
+ case OMP_CLAUSE_NOCONTEXT:
pc = &OMP_CLAUSE_CHAIN (c);
continue;
new file mode 100644
@@ -0,0 +1,32 @@
+/* Test parsing of OMP clause adjust_args */
+/* { dg-do compile } */
+
+int b;
+
+int f0 (void *a);
+int g (void *a);
+int f1 (int);
+
+#pragma omp declare variant (f0) match (construct={target}) adjust_args (nothing: a) /* { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause" } */
+int f2 (void *a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (other: a) /* { dg-error "expected 'nothing' or 'need_device_ptr'" } */
+int f3 (int a);
+#pragma omp declare variant (f0) adjust_args (nothing: a) /* { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause" } */
+int f4 (void *a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args () /* { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" } */
+int f5 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing) /* { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" } */
+int f6 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing:) /* { dg-error "expected identifier before '\\)' token" } */
+int f7 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing: z) /* { dg-error "'z' undeclared here \\(not in a function\\)" } */
+int f8 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (need_device_ptr: a) /* { dg-error "'a' is not a C pointer" } */
+int f9 (int a);
+#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (nothing: a) /* { dg-error "'a' is specified more than once" } */
+int f10 (int a);
+#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: a) /* { dg-error "'a' is specified more than once" } */
+int f11 (void *a);
+#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: b) /* { dg-error "'b' is not a function argument" } */
+int f12 (void *a);
+
new file mode 100644
@@ -0,0 +1,53 @@
+/* Test parsing of #pragma omp dispatch */
+/* { dg-do compile } */
+
+int f0 (int);
+
+void f1 (void)
+{
+ int a, b;
+ double x;
+ struct {int a; int b;} s;
+ int arr[1];
+
+#pragma omp dispatch
+ int c = f0 (a); /* { dg-error "expected expression before 'int'" } */
+#pragma omp dispatch
+ int f2 (int d); /* { dg-error "expected expression before 'int'" } */
+#pragma omp dispatch
+ a = b; /* { dg-error "expected '\\(' before ';' token" } */
+#pragma omp dispatch
+ s.a = f0(a) + b; /* { dg-error "expected ';' before '\\+' token" } */
+#pragma omp dispatch
+ b = !f0(a); /* { dg-error "expected a function name before '!' token" } */
+#pragma omp dispatch
+ s.b += f0(s.a); /* { dg-error "expected '=' before '\\+=' token" } */
+#pragma omp dispatch
+#pragma omp threadprivate(a) /* { dg-error "expected expression before '#pragma'" } */
+ a = f0(b);
+
+#pragma omp dispatch nocontext(s) /* { dg-error "used struct type value where scalar is required" } */
+ f0(a);
+#pragma omp dispatch nocontext(a, b) /* { dg-error "expected '\\)' before ','" } */
+ f0(a);
+#pragma omp dispatch nocontext(a) nocontext(b) /* { dg-error "too many 'nocontext' clauses" } */
+ f0(a);
+#pragma omp dispatch novariants(s) /* { dg-error "used struct type value where scalar is required" } */
+ f0(a);
+#pragma omp dispatch novariants(a, b) /* { dg-error "expected '\\)' before ','" } */
+ f0(a);
+#pragma omp dispatch novariants(a) novariants(b) /* { dg-error "too many 'novariants' clauses" } */
+ f0(a);
+#pragma omp dispatch nowait nowait /* { dg-error "too many 'nowait' clauses" } */
+ f0(a);
+#pragma omp dispatch device(x) /* { dg-error "expected integer expression before end of line" } */
+ f0(a);
+#pragma omp dispatch device(arr) /* { dg-error "expected integer expression before end of line" } */
+ f0(a);
+#pragma omp dispatch is_device_ptr(x) /* { dg-error "'is_device_ptr' variable is neither a pointer nor an array" } */
+ f0(a);
+#pragma omp dispatch is_device_ptr(&x) /* { dg-error "expected identifier before '&' token" } */
+ f0(a);
+#pragma omp dispatch depend(inout: f0) /* { dg-error "'f0' is not lvalue expression nor array section in 'depend' clause" } */
+ f0(a);
+}
new file mode 100644
@@ -0,0 +1,76 @@
+// Adapted from OpenMP examples
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+int baz (double *d_bv, const double *d_av, int n)
+{
+#pragma omp distribute parallel for
+ for (int i = 0; i < n; i++)
+ d_bv[i] = d_av[i] * i;
+ return -3;
+}
+
+int bar (double *d_bv, const double *d_av, int n)
+{
+#pragma omp target is_device_ptr(d_bv, d_av)
+ for (int i = 0; i < n; i++)
+ d_bv[i] = d_av[i] * i;
+ return -2;
+}
+
+#pragma omp declare variant(bar) match(construct={dispatch}) adjust_args(need_device_ptr: bv, av)
+#pragma omp declare variant(baz) match(implementation={vendor(gnu)})
+int foo (double *bv, const double *av, int n)
+{
+ for (int i = 0; i < n; i++)
+ bv[i] = av[i] * i;
+ return -1;
+}
+
+int test (int n)
+{
+ const double e = 2.71828;
+
+ double *av = (double *) malloc (n * sizeof (*av));
+ double *bv = (double *) malloc (n * sizeof (*bv));
+ double *d_bv = (double *) malloc (n * sizeof (*d_bv));
+
+ for (int i = 0; i < n; i++)
+ {
+ av[i] = e * i;
+ bv[i] = 0.0;
+ d_bv[i] = 0.0;
+ }
+
+ int f, last_dev = omp_get_num_devices () - 1;
+#pragma omp target data map(to: av[:n]) map(from: d_bv[:n]) device(last_dev) if (n == 1024)
+ {
+ #pragma omp dispatch nocontext(n > 1024) novariants(n < 1024) device(last_dev)
+ f = foo (d_bv, av, n);
+ }
+
+ foo (bv, av, n);
+ for (int i = 0; i < n; i++)
+ {
+ if (d_bv[i] != bv[i])
+ {
+ fprintf (stderr, "ERROR at %d: %lf (act) != %lf (exp)\n", i, d_bv[i], bv[i]);
+ return 1;
+ }
+ }
+ return f;
+}
+
+int
+main (void)
+{
+ int ret = test(1023);
+ if (ret != -1) return 1;
+ ret = test(1024);
+ if (ret != -2) return 1;
+ ret = test(1025);
+ if (ret != -3) return 1;
+ return 0;
+}