case OMP_CLAUSE_IS_DEVICE_PTR:
case OMP_CLAUSE_HAS_DEVICE_ADDR:
case OMP_CLAUSE_MAP:
+ case OMP_CLAUSE_USES_ALLOCATORS:
s = C_OMP_CLAUSE_SPLIT_TARGET;
break;
case OMP_CLAUSE_DOACROSS:
PRAGMA_OMP_CLAUSE_USE,
PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR,
PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR,
+ PRAGMA_OMP_CLAUSE_USES_ALLOCATORS,
/* Clauses for OpenACC. */
PRAGMA_OACC_CLAUSE_ASYNC,
result = PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR;
else if (!strcmp ("use_device_ptr", p))
result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR;
+ else if (!strcmp ("uses_allocators", p))
+ result = PRAGMA_OMP_CLAUSE_USES_ALLOCATORS;
break;
case 'v':
if (!strcmp ("vector", p))
return nl;
}
+/* OpenMP 5.0:
+ uses_allocators ( allocator-list )
+
+ allocator-list:
+ allocator
+ allocator , allocator-list
+ allocator ( traits-array )
+ allocator ( traits-array ) , allocator-list
+
+ OpenMP 5.2:
+
+ uses_allocators ( modifier : allocator-list )
+ uses_allocators ( modifier , modifier : allocator-list )
+
+ modifier:
+ traits ( traits-array )
+ memspace ( mem-space-handle ) */
+
+static tree
+c_parser_omp_clause_uses_allocators (c_parser *parser, tree list)
+{
+ location_t clause_loc = c_parser_peek_token (parser)->location;
+ tree nl = list;
+ matching_parens parens;
+ if (!parens.require_open (parser))
+ return list;
+
+ bool has_modifiers = false;
+ bool seen_allocators = false;
+ tree memspace_expr = NULL_TREE;
+ tree traits_var = NULL_TREE;
+
+ if (c_parser_next_token_is (parser, CPP_NAME)
+ && c_parser_peek_2nd_token (parser)->type == CPP_OPEN_PAREN)
+ {
+ unsigned int n = 3;
+ const char *p
+ = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+ if ((strcmp (p, "traits") == 0 || strcmp (p, "memspace") == 0)
+ && c_parser_check_balanced_raw_token_sequence (parser, &n)
+ && (c_parser_peek_nth_token_raw (parser, n)->type
+ == CPP_CLOSE_PAREN))
+ {
+ if (c_parser_peek_nth_token_raw (parser, n + 1)->type
+ == CPP_COLON)
+ has_modifiers = true;
+ else if (c_parser_peek_nth_token_raw (parser, n + 1)->type
+ == CPP_COMMA
+ && (c_parser_peek_nth_token_raw (parser, n + 2)->type
+ == CPP_NAME)
+ && (c_parser_peek_nth_token_raw (parser, n + 3)->type
+ == CPP_OPEN_PAREN))
+ {
+ c_token *tok = c_parser_peek_nth_token_raw (parser, n + 2);
+ const char *q = IDENTIFIER_POINTER (tok->value);
+ n += 4;
+ if ((strcmp (q, "traits") == 0
+ || strcmp (q, "memspace") == 0)
+ && c_parser_check_balanced_raw_token_sequence (parser, &n)
+ && (c_parser_peek_nth_token_raw (parser, n)->type
+ == CPP_CLOSE_PAREN))
+ {
+ if (c_parser_peek_nth_token_raw (parser, n + 1)->type
+ == CPP_COLON)
+ has_modifiers = true;
+ if ((c_parser_peek_nth_token_raw (parser, n + 1)->type
+ == CPP_COMMA)
+ && (c_parser_peek_nth_token_raw (parser, n + 2)->type
+ == CPP_NAME))
+ {
+ c_token *tok
+ = c_parser_peek_nth_token_raw (parser, n + 2);
+ const char *m = IDENTIFIER_POINTER (tok->value);
+ if (strcmp (p, m) == 0 || strcmp (q, m) == 0)
+ {
+ error_at (tok->location, "duplicate %qs modifier", m);
+ goto end;
+ }
+ }
+ }
+ }
+ }
+ if (has_modifiers)
+ {
+ c_parser_consume_token (parser);
+ matching_parens parens2;
+ parens2.require_open (parser);
+ c_expr expr = c_parser_expr_no_commas (parser, NULL);
+ if (expr.value == error_mark_node)
+ ;
+ else if (strcmp (p, "traits") == 0)
+ {
+ traits_var = expr.value;
+ traits_var = c_fully_fold (traits_var, false, NULL);
+ }
+ else
+ {
+ memspace_expr = expr.value;
+ memspace_expr = c_fully_fold (memspace_expr, false, NULL);
+ }
+ parens2.skip_until_found_close (parser);
+ if (c_parser_next_token_is (parser, CPP_COMMA))
+ {
+ c_parser_consume_token (parser);
+ c_token *tok = c_parser_peek_token (parser);
+ const char *q = "";
+ if (c_parser_next_token_is (parser, CPP_NAME))
+ q = IDENTIFIER_POINTER (tok->value);
+ if (strcmp (q, "traits") != 0 && strcmp (q, "memspace") != 0)
+ {
+ c_parser_error (parser, "expected %<traits%> or "
+ "%<memspace%>");
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+ else if (strcmp (p, q) == 0)
+ {
+ error_at (tok->location, "duplicate %qs modifier", p);
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+ c_parser_consume_token (parser);
+ if (!parens2.require_open (parser))
+ {
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+ expr = c_parser_expr_no_commas (parser, NULL);
+ if (strcmp (q, "traits") == 0)
+ {
+ traits_var = expr.value;
+ traits_var = c_fully_fold (traits_var, false, NULL);
+ }
+ else
+ {
+ memspace_expr = expr.value;
+ memspace_expr = c_fully_fold (memspace_expr, false, NULL);
+ }
+ parens2.skip_until_found_close (parser);
+ }
+ if (!c_parser_require (parser, CPP_COLON, "expected %<:%>"))
+ goto end;
+ }
+ }
+
+ while (c_parser_next_token_is (parser, CPP_NAME))
+ {
+ location_t alloc_loc = c_parser_peek_token (parser)->location;
+ c_token *tok = c_parser_peek_token (parser);
+ const char *tok_s = IDENTIFIER_POINTER (tok->value);
+ tree t = lookup_name (tok->value);
+ if (t == NULL_TREE)
+ {
+ undeclared_variable (tok->location, tok->value);
+ t = error_mark_node;
+ }
+ c_parser_consume_token (parser);
+
+ /* Legacy traits syntax. */
+ tree legacy_traits = NULL_TREE;
+ if (c_parser_next_token_is (parser, CPP_OPEN_PAREN)
+ && c_parser_peek_2nd_token (parser)->type == CPP_NAME
+ && c_parser_peek_nth_token_raw (parser, 3)->type == CPP_CLOSE_PAREN)
+ {
+ matching_parens parens2;
+ parens2.require_open (parser);
+ const char *tok_a
+ = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+ c_expr expr = c_parser_expr_no_commas (parser, NULL);
+ location_t close_loc = c_parser_peek_token (parser)->location;
+ parens2.skip_until_found_close (parser);
+
+ if (has_modifiers)
+ {
+ error_at (make_location (alloc_loc, alloc_loc, close_loc),
+ "legacy %<%s(%s)%> traits syntax not allowed in "
+ "%<uses_allocators%> clause when using modifiers",
+ tok_s, tok_a);
+ goto end;
+ }
+ legacy_traits = c_fully_fold (expr.value, false, NULL);
+ if (legacy_traits == error_mark_node)
+ goto end;
+
+ gcc_rich_location richloc (make_location (alloc_loc, alloc_loc, close_loc));
+ if (nl == list)
+ {
+ /* Fixit only works well if it is the only first item. */
+ richloc.add_fixit_replace (alloc_loc, "traits");
+ richloc.add_fixit_insert_after (close_loc, ": ");
+ richloc.add_fixit_insert_after (close_loc, tok_s);
+ }
+ warning_at (&richloc, OPT_Wdeprecated_openmp,
+ "the specification of arguments to %<uses_allocators%> "
+ "where each item is of the form %<allocator(traits)%> is "
+ "deprecated since OpenMP 5.2");
+ }
+
+ if (seen_allocators && has_modifiers)
+ {
+ error_at (c_parser_peek_token (parser)->location,
+ "%<uses_allocators%> clause only accepts a single "
+ "allocator when using modifiers");
+ goto end;
+ }
+ seen_allocators = true;
+
+ tree c = build_omp_clause (clause_loc,
+ OMP_CLAUSE_USES_ALLOCATORS);
+ OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c) = t;
+ OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c) = memspace_expr;
+ OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c) = (legacy_traits
+ ? legacy_traits : traits_var);
+ OMP_CLAUSE_CHAIN (c) = nl;
+ nl = c;
+
+ if (c_parser_next_token_is (parser, CPP_COMMA))
+ c_parser_consume_token (parser);
+ else
+ break;
+ }
+
+ end:
+ parens.skip_until_found_close (parser);
+ return nl;
+}
+
/* OpenMP 4.0:
linear ( variable-list )
linear ( variable-list : expression )
clauses = c_parser_omp_clause_linear (parser, clauses);
c_name = "linear";
break;
+ case PRAGMA_OMP_CLAUSE_USES_ALLOCATORS:
+ clauses = c_parser_omp_clause_uses_allocators (parser, clauses);
+ c_name = "uses_allocators";
+ break;
case PRAGMA_OMP_CLAUSE_AFFINITY:
clauses = c_parser_omp_clause_affinity (parser, clauses);
c_name = "affinity";
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_THREAD_LIMIT))
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_THREAD_LIMIT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USES_ALLOCATORS))
static bool
c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p)
break;
}
gcc_unreachable ();
+
+ case OMP_CLAUSE_USES_ALLOCATORS:
+ t = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);
+ if (t == error_mark_node)
+ {
+ remove = true;
+ break;
+ }
+ if ((VAR_P (t) || TREE_CODE (t) == PARM_DECL)
+ && (bitmap_bit_p (&generic_head, DECL_UID (t))
+ || bitmap_bit_p (&map_head, DECL_UID (t))
+ || bitmap_bit_p (&firstprivate_head, DECL_UID (t))
+ || bitmap_bit_p (&lastprivate_head, DECL_UID (t))))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qE appears more than once in data clauses", t);
+ remove = true;
+ break;
+ }
+ else
+ bitmap_set_bit (&generic_head, DECL_UID (t));
+ if (TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE
+ || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE (t))),
+ "omp_allocator_handle_t") != 0)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "allocator %qE must be of %<omp_allocator_handle_t%> "
+ "type", t);
+ remove = true;
+ break;
+ }
+ tree init;
+ if (!DECL_P (t)
+ || (TREE_CODE (t) == CONST_DECL
+ && ((init = DECL_INITIAL(t)) == nullptr
+ || TREE_CODE (init) != INTEGER_CST
+ || ((wi::to_widest (init) < 0
+ || wi::to_widest (init) > GOMP_OMP_PREDEF_ALLOC_MAX)
+ && (wi::to_widest (init) < GOMP_OMPX_PREDEF_ALLOC_MIN
+ || (wi::to_widest (init)
+ > GOMP_OMPX_PREDEF_ALLOC_MAX))))))
+ {
+ remove = true;
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "allocator %qE must be either a variable or a "
+ "predefined allocator", t);
+ break;
+ }
+ else if (TREE_CODE (t) == CONST_DECL)
+ {
+ /* omp_null_allocator is ignored and for predefined allocators,
+ not special handling is required; thus, remove them removed. */
+ remove = true;
+
+ if (OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c)
+ || OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "modifiers cannot be used with predefined "
+ "allocator %qE", t);
+ break;
+ }
+ }
+ t = OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c);
+ if (t == error_mark_node)
+ {
+ remove = true;
+ break;
+ }
+ if (t != NULL_TREE
+ && ((TREE_CODE (t) != CONST_DECL && TREE_CODE (t) != INTEGER_CST)
+ || TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE
+ || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE (t))),
+ "omp_memspace_handle_t") != 0))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "memspace modifier %qE must be"
+ " constant enum of %<omp_memspace_handle_t%> type", t);
+ remove = true;
+ break;
+ }
+ t = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c);
+ if (t == error_mark_node)
+ {
+ remove = true;
+ break;
+ }
+ if (t != NULL_TREE
+ && t != error_mark_node
+ && (DECL_EXTERNAL (t)
+ || TREE_CODE (t) == PARM_DECL))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "traits array %qE must be "
+ "defined in same scope as the construct on which the "
+ "clause appears", t);
+ remove = true;
+ }
+ if (t != NULL_TREE)
+ {
+ bool type_err = false;
+
+ if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE
+ || DECL_SIZE (t) == NULL_TREE
+ || !COMPLETE_TYPE_P (TREE_TYPE (t)))
+ type_err = true;
+ else
+ {
+ tree elem_t = TREE_TYPE (TREE_TYPE (t));
+ if (TREE_CODE (elem_t) != RECORD_TYPE
+ || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (elem_t)),
+ "omp_alloctrait_t") != 0
+ || !TYPE_READONLY (elem_t))
+ type_err = true;
+ }
+ if (type_err)
+ {
+ if (t != error_mark_node)
+ error_at (OMP_CLAUSE_LOCATION (c), "traits array %qE must "
+ "be of %<const omp_alloctrait_t []%> type", t);
+ else
+ error_at (OMP_CLAUSE_LOCATION (c), "traits array must "
+ "be of %<const omp_alloctrait_t []%> type");
+ remove = true;
+ }
+ else
+ {
+ tree cst_val = decl_constant_value_1 (t, true);
+ if (cst_val == t)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "traits array must be "
+ "initialized with constants");
+
+ remove = true;
+ }
+ }
+ }
+ if (remove)
+ break;
+ pc = &OMP_CLAUSE_CHAIN (c);
+ continue;
case OMP_CLAUSE_DEPEND:
depend_clause = c;
/* FALLTHRU */
result = PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR;
else if (!strcmp ("use_device_ptr", p))
result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR;
+ else if (!strcmp ("uses_allocators", p))
+ result = PRAGMA_OMP_CLAUSE_USES_ALLOCATORS;
break;
case 'v':
if (!strcmp ("vector", p))
return nlist;
}
+/* OpenMP 5.0:
+ uses_allocators ( allocator-list )
+
+ allocator-list:
+ allocator
+ allocator , allocator-list
+ allocator ( traits-array )
+ allocator ( traits-array ) , allocator-list
+
+ OpenMP 5.2:
+
+ uses_allocators ( modifier : allocator-list )
+ uses_allocators ( modifier , modifier : allocator-list )
+
+ modifier:
+ traits ( traits-array )
+ memspace ( mem-space-handle ) */
+
+static tree
+cp_parser_omp_clause_uses_allocators (cp_parser *parser, tree list)
+{
+ location_t clause_loc
+ = cp_lexer_peek_token (parser->lexer)->location;
+ tree nl = list;
+ matching_parens parens;
+ if (!parens.require_open (parser))
+ return list;
+
+ bool has_modifiers = false;
+ bool seen_allocators = false;
+ tree memspace_expr = NULL_TREE;
+ tree traits_var = NULL_TREE;
+
+ cp_parser_parse_tentatively (parser);
+ bool saved_colon_corrects_to_scope_p = parser->colon_corrects_to_scope_p;
+ parser->colon_corrects_to_scope_p = false;
+
+ cp_token *dup_mod_tok = NULL;
+ for (int mod = 0; mod <= 2; mod++)
+ if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)
+ && cp_lexer_nth_token_is (parser->lexer, 2, CPP_OPEN_PAREN))
+ {
+ cp_token *mod_tok = cp_lexer_peek_token (parser->lexer);
+ tree id = mod_tok->u.value;
+ const char *p = IDENTIFIER_POINTER (id);
+ if (strcmp (p, "traits") != 0 && strcmp (p, "memspace") != 0)
+ break;
+ cp_lexer_consume_token (parser->lexer);
+ matching_parens parens2;
+ if (!parens2.require_open (parser))
+ break;
+ tree t = cp_parser_assignment_expression (parser);
+ if (strcmp (p, "traits") == 0)
+ {
+ if (traits_var != NULL_TREE)
+ dup_mod_tok = mod_tok;
+ else
+ traits_var = t;
+ }
+ else
+ {
+ if (memspace_expr != NULL_TREE)
+ dup_mod_tok = mod_tok;
+ else
+ memspace_expr = t;
+ }
+ if (!parens2.require_close (parser))
+ break;
+ if (cp_lexer_next_token_is (parser->lexer, CPP_COLON))
+ {
+ has_modifiers = true;
+ cp_lexer_consume_token (parser->lexer);
+ break;
+ }
+ if (/*mod != 0 || */ cp_lexer_next_token_is_not (parser->lexer, CPP_COMMA))
+ break;
+ cp_lexer_consume_token (parser->lexer);
+ }
+ else
+ break;
+
+ if (!has_modifiers)
+ {
+ cp_parser_abort_tentative_parse (parser);
+ traits_var = NULL_TREE;
+ memspace_expr = NULL_TREE;
+ }
+ else
+ {
+ if (dup_mod_tok)
+ {
+ error_at (dup_mod_tok->location, "duplicate %qs modifier",
+ IDENTIFIER_POINTER (dup_mod_tok->u.value));
+ cp_parser_parse_definitely (parser);
+ goto end;
+ }
+ cp_parser_parse_definitely (parser);
+ }
+ parser->colon_corrects_to_scope_p = saved_colon_corrects_to_scope_p;
+
+ while (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+ {
+ cp_token *tok = cp_lexer_peek_token (parser->lexer);
+ tree t;
+ t = cp_parser_lookup_name_simple (parser,
+ tok->u.value,
+ tok->location);
+ if (t == error_mark_node)
+ cp_parser_name_lookup_error (parser, tok->u.value, t, NLE_NULL,
+ tok->location);
+ cp_lexer_consume_token (parser->lexer);
+
+ /* Legacy traits syntax. */
+ tree legacy_traits = NULL_TREE;
+ if (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_PAREN)
+ && cp_lexer_nth_token_is (parser->lexer, 2, CPP_NAME)
+ && cp_lexer_nth_token_is (parser->lexer, 3, CPP_CLOSE_PAREN))
+ {
+ matching_parens parens2;
+ parens2.require_open (parser);
+ cp_token *arg_tok = cp_lexer_peek_token (parser->lexer);
+
+ tree arg = cp_parser_lookup_name_simple (parser, arg_tok->u.value,
+ arg_tok->location);
+ if (arg == error_mark_node)
+ cp_parser_name_lookup_error (parser, arg_tok->u.value, arg,
+ NLE_NULL, arg_tok->location);
+ cp_lexer_consume_token (parser->lexer);
+ location_t close_loc = cp_lexer_peek_token (parser->lexer)->location;
+ parens2.require_close (parser);
+
+ if (has_modifiers)
+ {
+ error_at (make_location (tok->location, tok->location, close_loc),
+ "legacy %<%E(%E)%> traits syntax not allowed in "
+ "%<uses_allocators%> clause when using modifiers",
+ tok->u.value, arg_tok->u.value);
+ goto end;
+ }
+ legacy_traits = arg;
+ if (legacy_traits == error_mark_node)
+ goto end;
+ gcc_rich_location richloc (make_location (tok->location,
+ tok->location, close_loc));
+ if (nl == list)
+ {
+ /* Fixit only works well if it is the first item. */
+ richloc.add_fixit_replace (tok->location, "traits");
+ richloc.add_fixit_insert_after (close_loc, ": ");
+ richloc.add_fixit_insert_after (close_loc,
+ IDENTIFIER_POINTER (tok->u.value));
+ }
+ warning_at (&richloc, OPT_Wdeprecated_openmp,
+ "the specification of arguments to %<uses_allocators%> "
+ "where each item is of the form %<allocator(traits)%> is "
+ "deprecated since OpenMP 5.2");
+ }
+
+ if (seen_allocators && has_modifiers)
+ {
+ error_at (cp_lexer_peek_token (parser->lexer)->location,
+ "%<uses_allocators%> clause only accepts a single "
+ "allocator when using modifiers");
+ goto end;
+ }
+ seen_allocators = true;
+
+ tree c = build_omp_clause (clause_loc,
+ OMP_CLAUSE_USES_ALLOCATORS);
+ OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c) = t;
+ OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c) = memspace_expr;
+ OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c) = (legacy_traits
+ ? legacy_traits : traits_var);
+ OMP_CLAUSE_CHAIN (c) = nl;
+ nl = c;
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
+ cp_lexer_consume_token (parser->lexer);
+ else
+ break;
+ }
+
+ if (!parens.require_close (parser))
+ goto end;
+ return nl;
+ end:
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return nl;
+}
+
/* OpenMP 2.5:
lastprivate ( variable-list )
clauses = cp_parser_omp_clause_allocate (parser, clauses);
c_name = "allocate";
break;
+ case PRAGMA_OMP_CLAUSE_USES_ALLOCATORS:
+ clauses = cp_parser_omp_clause_uses_allocators (parser, clauses);
+ c_name = "uses_allocators";
+ break;
case PRAGMA_OMP_CLAUSE_LINEAR:
{
bool declare_simd = false;
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_THREAD_LIMIT))
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_THREAD_LIMIT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USES_ALLOCATORS))
static bool
cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
= tsubst_stmt (OMP_CLAUSE_LINEAR_STEP (oc), args,
complain, in_decl);
break;
+ case OMP_CLAUSE_USES_ALLOCATORS:
+ OMP_CLAUSE_OPERAND (nc, 0)
+ = tsubst_stmt (OMP_CLAUSE_OPERAND (oc, 0), args, complain, in_decl);
+ OMP_CLAUSE_OPERAND (nc, 1)
+ = tsubst_stmt (OMP_CLAUSE_OPERAND (oc, 1), args, complain, in_decl);
+ OMP_CLAUSE_OPERAND (nc, 2)
+ = tsubst_stmt (OMP_CLAUSE_OPERAND (oc, 2), args, complain, in_decl);
+ break;
case OMP_CLAUSE_INIT:
if (ort == C_ORT_OMP_INTEROP
&& OMP_CLAUSE_INIT_PREFER_TYPE (nc)
break;
}
gcc_unreachable ();
+ case OMP_CLAUSE_USES_ALLOCATORS:
+ t = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);
+ t = convert_from_reference (t);
+ if (t == error_mark_node)
+ {
+ remove = true;
+ break;
+ }
+ if (DECL_P (t)
+ && (bitmap_bit_p (&generic_head, DECL_UID (t))
+ || bitmap_bit_p (&firstprivate_head, DECL_UID (t))
+ || bitmap_bit_p (&map_firstprivate_head, DECL_UID (t))))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qE appears more than once in data clauses", t);
+ remove = true;
+ }
+ else if (DECL_P (t))
+ bitmap_set_bit (&generic_head, DECL_UID (t));
+ if (type_dependent_expression_p (t))
+ break;
+ if (TREE_CODE (t) == FIELD_DECL)
+ {
+ sorry_at (OMP_CLAUSE_LOCATION (c), "class member %qE not yet "
+ "supported in %<uses_allocators%> clause", t);
+ remove = true;
+ break;
+ }
+ if (TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE
+ || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE (t))),
+ "omp_allocator_handle_t") != 0)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "allocator %qE must be of %<omp_allocator_handle_t%> "
+ "type", t);
+ remove = true;
+ break;
+ }
+ tree init;
+ if (TREE_CODE (t) == CONST_DECL)
+ init = DECL_INITIAL(t);
+ else
+ init = t;
+ if (!DECL_P (t)
+ && (init == NULL_TREE
+ || TREE_CODE (init) != INTEGER_CST
+ || ((wi::to_widest (init) < 0
+ || wi::to_widest (init) > GOMP_OMP_PREDEF_ALLOC_MAX)
+ && (wi::to_widest (init) < GOMP_OMPX_PREDEF_ALLOC_MIN
+ || (wi::to_widest (init)
+ > GOMP_OMPX_PREDEF_ALLOC_MAX)))))
+ {
+ remove = true;
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "allocator %qE must be either a variable or a "
+ "predefined allocator", t);
+ break;
+ }
+ else if (TREE_CODE (t) == CONST_DECL)
+ {
+ /* omp_null_allocator is ignored and for predefined allocators,
+ not special handling is required; thus, remove them removed. */
+ remove = true;
+
+ if (OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c)
+ || OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "modifiers cannot be used with predefined "
+ "allocators");
+ break;
+ }
+ }
+ t = OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c);
+ if (t == error_mark_node)
+ {
+ remove = true;
+ break;
+ }
+ if (t != NULL_TREE
+ && !type_dependent_expression_p (t)
+ && ((TREE_CODE (t) != CONST_DECL && TREE_CODE (t) != INTEGER_CST)
+ || TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE
+ || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE (t))),
+ "omp_memspace_handle_t") != 0))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "memspace modifier %qE must be"
+ " constant enum of %<omp_memspace_handle_t%> type", t);
+ remove = true;
+ break;
+ }
+ t = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c);
+ if (t == error_mark_node)
+ {
+ remove = true;
+ break;
+ }
+ if (type_dependent_expression_p (t))
+ break;
+ if (t != NULL_TREE
+ && t != error_mark_node
+ && !type_dependent_expression_p (t)
+ && (!DECL_P (t)
+ || DECL_EXTERNAL (t)
+ || TREE_CODE (t) == PARM_DECL))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "traits array %qE must be "
+ "defined in same scope as the construct on which the "
+ "clause appears", t);
+ remove = true;
+ }
+ if (t != NULL_TREE)
+ {
+ bool type_err = false;
+
+ if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE
+ || DECL_SIZE (t) == NULL_TREE
+ || !COMPLETE_TYPE_P (TREE_TYPE (t)))
+ type_err = true;
+ else
+ {
+ tree elem_t = TREE_TYPE (TREE_TYPE (t));
+ if (TREE_CODE (elem_t) != RECORD_TYPE
+ || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (elem_t)),
+ "omp_alloctrait_t") != 0
+ || !TYPE_READONLY (elem_t))
+ type_err = true;
+ }
+ if (type_err)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "traits array %qE must "
+ "be of %<const omp_alloctrait_t []%> type", t);
+ remove = true;
+ }
+ else if (TREE_CODE (array_type_nelts_top (TREE_TYPE (t)))
+ != INTEGER_CST)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "variable length traits "
+ "arrays are not supported");
+ remove = true;
+ }
+ else
+ {
+ tree cst_val = decl_constant_value (t);
+ if (cst_val == t)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "traits array must be "
+ "initialized with constants");
+ remove = true;
+ }
+ }
+ }
+ if (remove)
+ break;
+ pc = &OMP_CLAUSE_CHAIN (c);
+ continue;
case OMP_CLAUSE_DEPEND:
depend_clause = c;
/* FALLTHRU */
"%<omp_allocator_handle_kind%>", n->sym->name,
&n->where);
else if (n->sym->attr.flavor != FL_VARIABLE
+ && strcmp (n->sym->name, "omp_null_allocator") != 0
&& ((!startswith (n->sym->name, "omp_")
&& !startswith (n->sym->name, "ompx_"))
|| !endswith (n->sym->name, "_mem_alloc")))
}
break;
case OMP_LIST_USES_ALLOCATORS:
- /* Ignore pre-defined allocators as no special treatment is needed. */
+ /* Ignore omp_null_allocator and pre-defined allocators as no
+ special treatment is needed. */
for (; n != NULL; n = n->next)
if (n->sym->attr.flavor == FL_VARIABLE)
break;
nowait = 1;
break;
+ case OMP_CLAUSE_USES_ALLOCATORS:
+ sorry_at (OMP_CLAUSE_LOCATION (c), "%<uses_allocators%> clause");
+ remove = 1;
+ break;
+
case OMP_CLAUSE_ORDERED:
case OMP_CLAUSE_UNTIED:
case OMP_CLAUSE_COLLAPSE:
case OMP_CLAUSE_FINALIZE:
case OMP_CLAUSE_INCLUSIVE:
case OMP_CLAUSE_EXCLUSIVE:
+ case OMP_CLAUSE_USES_ALLOCATORS:
break;
case OMP_CLAUSE_NOHOST:
--- /dev/null
+typedef enum omp_allocator_handle_t
+#if __cplusplus >= 201103L
+: __UINTPTR_TYPE__
+#endif
+{
+ omp_default_mem_alloc = 1,
+ omp_low_lat_mem_alloc = 5,
+ __omp_allocator_handle_t_max__ = __UINTPTR_MAX__
+} omp_allocator_handle_t;
+
+typedef struct omp_alloctrait_t
+{
+ int key;
+ int value;
+} omp_alloctrait_t;
+
+extern void *omp_alloc (__SIZE_TYPE__, omp_allocator_handle_t);
+
+void
+f (omp_allocator_handle_t my_alloc)
+{
+ #pragma omp target
+ {
+ int a; /* { dg-error "'my_alloc' in 'allocator' clause inside a target region must be specified in an 'uses_allocators' clause on the 'target' directive" "not yet implemented" { xfail *-*-* } } */
+ #pragma omp allocate(a) allocator(my_alloc) /* { dg-message "sorry, unimplemented: '#pragma omp allocate' not yet supported" "" { target c++ } } */
+ a = 5;
+ void *prt = omp_alloc(32, my_alloc);
+ #pragma omp parallel allocate(allocator(my_alloc) : a) firstprivate(a) /* { dg-error "allocator 'my_alloc' in 'allocate' clause inside a target region must be specified in an 'uses_allocators' clause on the 'target' directive" "not yet implemented" { xfail *-*-* } } */
+ a = 7;
+ }
+}
+
+void
+g (omp_allocator_handle_t my_alloc)
+{
+ /* The following defines a default-mem-space allocator with no extra traits. */
+ #pragma omp target uses_allocators(my_alloc)
+ {
+ int a;
+ #pragma omp allocate(a) allocator(my_alloc) /* { dg-message "sorry, unimplemented: '#pragma omp allocate' not yet supported" "" { target c++ } } */
+ a = 5;
+ void *prt = omp_alloc(32, my_alloc);
+ #pragma omp parallel allocate(allocator(my_alloc) : a) firstprivate(a)
+ a = 7;
+ }
+}
+
+// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 37 }
--- /dev/null
+typedef enum omp_allocator_handle_t
+#if __cplusplus >= 201103L
+: __UINTPTR_TYPE__
+#endif
+{
+ omp_default_mem_alloc = 1,
+ omp_low_lat_mem_alloc = 5,
+ __omp_allocator_handle_t_max__ = __UINTPTR_MAX__
+} omp_allocator_handle_t;
+
+typedef struct omp_alloctrait_t
+{
+ int key;
+ int value;
+} omp_alloctrait_t;
+
+void
+f ()
+{
+ omp_alloctrait_t trait[1] = {{1,1}};
+ omp_allocator_handle_t my_alloc;
+ #pragma omp target uses_allocators(traits(trait) : my_alloc) /* { dg-error "traits array 'trait' must be of 'const omp_alloctrait_t \\\[\\\]' type" } */
+ ;
+}
+
+void
+g ()
+{
+ const omp_alloctrait_t trait[1] = {{1,1}};
+ omp_allocator_handle_t my_alloc;
+ #pragma omp target uses_allocators(traits(trait) : my_alloc)
+ ;
+}
+
+// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 31 }
--- /dev/null
+/* { dg-do compile } */
+/* { dg-additional-options "-Wno-deprecated-openmp" } */
+
+//#include <omp.h>
+
+typedef __UINTPTR_TYPE__ omp_uintptr_t;
+
+#if __cplusplus >= 201103L
+# define __GOMP_UINTPTR_T_ENUM : omp_uintptr_t
+#else
+# define __GOMP_UINTPTR_T_ENUM
+#endif
+
+typedef enum omp_memspace_handle_t __GOMP_UINTPTR_T_ENUM
+{
+ omp_default_mem_space = 0,
+ omp_large_cap_mem_space = 1,
+ omp_const_mem_space = 2,
+ omp_high_bw_mem_space = 3,
+ omp_low_lat_mem_space = 4,
+ ompx_gnu_managed_mem_space = 200,
+ __omp_memspace_handle_t_max__ = __UINTPTR_MAX__
+} omp_memspace_handle_t;
+
+typedef enum omp_allocator_handle_t __GOMP_UINTPTR_T_ENUM
+{
+ omp_null_allocator = 0,
+ omp_default_mem_alloc = 1,
+ omp_large_cap_mem_alloc = 2,
+ omp_const_mem_alloc = 3,
+ omp_high_bw_mem_alloc = 4,
+ omp_low_lat_mem_alloc = 5,
+ omp_cgroup_mem_alloc = 6,
+ omp_pteam_mem_alloc = 7,
+ omp_thread_mem_alloc = 8,
+ ompx_gnu_pinned_mem_alloc = 200,
+ ompx_gnu_managed_mem_alloc = 201,
+ __omp_allocator_handle_t_max__ = __UINTPTR_MAX__
+} omp_allocator_handle_t;
+
+typedef enum omp_alloctrait_key_t
+{
+ omp_atk_sync_hint = 1,
+ omp_atk_alignment = 2,
+ omp_atk_access = 3,
+ omp_atk_pool_size = 4,
+ omp_atk_fallback = 5,
+ omp_atk_fb_data = 6,
+ omp_atk_pinned = 7,
+ omp_atk_partition = 8
+} omp_alloctrait_key_t;
+
+typedef enum omp_alloctrait_value_t
+{
+ omp_atv_default = (__UINTPTR_TYPE__) -1,
+ omp_atv_false = 0,
+ omp_atv_true = 1,
+ omp_atv_contended = 3,
+ omp_atv_uncontended = 4,
+ omp_atv_serialized = 5,
+ omp_atv_private = 6,
+ omp_atv_all = 7,
+ omp_atv_thread = 8,
+ omp_atv_pteam = 9,
+ omp_atv_cgroup = 10,
+ omp_atv_default_mem_fb = 11,
+ omp_atv_null_fb = 12,
+ omp_atv_abort_fb = 13,
+ omp_atv_allocator_fb = 14,
+ omp_atv_environment = 15,
+ omp_atv_nearest = 16,
+ omp_atv_blocked = 17,
+ omp_atv_interleaved = 18
+} omp_alloctrait_value_t;
+
+typedef struct omp_alloctrait_t
+{
+ omp_alloctrait_key_t key;
+ omp_uintptr_t value;
+} omp_alloctrait_t;
+
+omp_alloctrait_key_t k;
+omp_alloctrait_value_t v;
+
+int f (const omp_alloctrait_t arg_traits[], int n)
+{
+ omp_allocator_handle_t foo, bar;
+ const omp_alloctrait_t traits_array[] = { { omp_atk_pinned, omp_atv_true },
+ { omp_atk_partition, omp_atv_nearest } };
+ extern const omp_alloctrait_t ex_traits[2];
+ extern const omp_alloctrait_t ex2_traits[];
+#ifndef __cplusplus
+ const omp_alloctrait_t vla_traits[n] = {}; /* Not useful, but shouldn't crash. */
+#else
+ const omp_alloctrait_t vla_traits[n] = { { omp_atk_pinned, omp_atv_true },
+ { omp_atk_partition, omp_atv_nearest } };
+#endif
+
+ #pragma omp target uses_allocators (baz) /* { dg-error "'baz' undeclared .first use in this function." "" { target c } } */
+ ; /* { dg-error "'baz' has not been declared" "" { target c++ } .-1 } */
+ #pragma omp target uses_allocators (foo (xyz)) /* { dg-error "'xyz' undeclared .first use in this function." "" { target c } } */
+ ; /* { dg-error "'xyz' has not been declared" "" { target c++ } .-1 } */
+ #pragma omp target uses_allocators (foo (traits_array), baz (traits_array)) /* { dg-error "'baz' has not been declared" "" { target c++ } } */
+ ;
+ #pragma omp target uses_allocators (foo (arg_traits)) /* { dg-error "traits array 'arg_traits' must be defined in same scope as the construct on which the clause appears" } */
+ ; /* { dg-error "traits array 'arg_traits' must be of 'const omp_alloctrait_t \\\[\\\]' type" "" { target *-*-* } .-1 } */
+ #pragma omp target uses_allocators (foo (ex_traits)) /* { dg-error "traits array 'ex_traits' must be defined in same scope as the construct on which the clause appears" } */
+ ; /* { dg-error "traits array must be initialized with constants" "" { target *-*-* } .-1 } */
+ #pragma omp target uses_allocators (foo (ex2_traits)) /* { dg-error "traits array 'ex2_traits' must be defined in same scope as the construct on which the clause appears" } */
+ ; /* { dg-error "traits array 'ex2_traits' must be of 'const omp_alloctrait_t \\\[\\\]' type" "" { target *-*-* } .-1 } */
+ #pragma omp target uses_allocators (foo (vla_traits)) /* { dg-error "variable length traits arrays are not supported" "" { target c++ } } */
+ ;
+ #pragma omp target uses_allocators (memspace(omp_no_such_space) : foo) /* { dg-error "'omp_no_such_space' undeclared .first use in this function." "" { target c } } */
+ ; /* { dg-error "'omp_no_such_space' was not declared in this scope" "" { target c++ } .-1 } */
+ #pragma omp target uses_allocators (memspace(1) : foo) /* { dg-error "memspace modifier '1' must be constant enum of 'omp_memspace_handle_t' type" } */
+ ;
+ #pragma omp target uses_allocators (memspace(omp_no_such_space) : foo, bar) /* { dg-error "'uses_allocators' clause only accepts a single allocator when using modifiers" } */
+ ; /* { dg-error "memspace modifier 'omp_no_such_space' must be constant enum of 'omp_memspace_handle_t' type" "" { target c++ } .-1 } */
+ #pragma omp target uses_allocators (traits(xyz) : bar) /* { dg-error "'xyz' was not declared in this scope" "" { target c++ } } */
+ ;
+ #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space), traits(traits_array), memspace (omp_no_such_space) : bar) /* { dg-error "duplicate 'memspace' modifier" "" { target c } } */
+ ; /* { dg-error "expected '\\\)' before 'memspace" "" { target c } .-1 } */
+ /* { dg-error "duplicate 'memspace' modifier" "" { target c++ } .-2 } */
+ #pragma omp target uses_allocators (traitz(traits_array), memspace(omp_high_bw_mem_space) : bar) /* { dg-error "'traitz' undeclared .first use in this function." "" { target c } } */
+ ; /* { dg-error "'memspace' undeclared .first use in this function." "" { target c } .-1 } */
+ /* { dg-error "'traitz' has not been declared" "" { target c++ } .-2 } */
+ /* { dg-error "'memspace' has not been declared" "" { target c++ } .-3 } */
+ /* { dg-error "expected '\\\)' before ':' token" "" { target *-*-* } .-4 } */
+ #pragma omp target uses_allocators (omp_null_allocator)
+ ;
+ #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo, bar) /* { dg-error "'uses_allocators' clause only accepts a single allocator when using modifiers" } */
+ ;
+ #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo(foo_traits)) /* { dg-error "'foo_traits' undeclared .first use in this function.; did you mean 'vla_traits'." "" { target c } } */
+ ; /* { dg-error "'foo_traits' has not been declared" "" { target c++ } .-1 } */
+ /* { dg-error "legacy 'foo\\\(foo_traits\\\)' traits syntax not allowed in 'uses_allocators' clause when using modifiers" "" { target *-*-* } .-2 } */
+ return 0;
+}
+
+// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 103 }
+// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target c } 111 }
+// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target c } 113 }
+// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target c } 117 }
+// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target c } 119 }
+// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 131 }
--- /dev/null
+// { dg-do compile }
+
+//#include <omp.h>
+
+typedef __UINTPTR_TYPE__ omp_uintptr_t;
+
+#if __cplusplus >= 201103L
+# define __GOMP_UINTPTR_T_ENUM : omp_uintptr_t
+#else
+# define __GOMP_UINTPTR_T_ENUM
+#endif
+
+typedef enum omp_memspace_handle_t __GOMP_UINTPTR_T_ENUM
+{
+ omp_default_mem_space = 0,
+ omp_large_cap_mem_space = 1,
+ omp_const_mem_space = 2,
+ omp_high_bw_mem_space = 3,
+ omp_low_lat_mem_space = 4,
+ ompx_gnu_managed_mem_space = 200,
+ __omp_memspace_handle_t_max__ = __UINTPTR_MAX__
+} omp_memspace_handle_t;
+
+typedef enum omp_allocator_handle_t __GOMP_UINTPTR_T_ENUM
+{
+ omp_null_allocator = 0,
+ omp_default_mem_alloc = 1,
+ omp_large_cap_mem_alloc = 2,
+ omp_const_mem_alloc = 3,
+ omp_high_bw_mem_alloc = 4,
+ omp_low_lat_mem_alloc = 5,
+ omp_cgroup_mem_alloc = 6,
+ omp_pteam_mem_alloc = 7,
+ omp_thread_mem_alloc = 8,
+ ompx_gnu_pinned_mem_alloc = 200,
+ ompx_gnu_managed_mem_alloc = 201,
+ __omp_allocator_handle_t_max__ = __UINTPTR_MAX__
+} omp_allocator_handle_t;
+
+typedef struct omp_alloctrait_t
+{
+// omp_alloctrait_key_t key;
+// omp_uintptr_t value;
+} omp_alloctrait_t;
+
+
+void f()
+{
+ omp_allocator_handle_t my;
+ struct t {
+ omp_allocator_handle_t a1;
+ } s;
+
+const omp_allocator_handle_t my3 = (omp_allocator_handle_t) 300;
+omp_allocator_handle_t my4[1];
+const omp_alloctrait_t t[] = {};
+ #pragma omp target uses_allocators(my, omp_default_mem_alloc, omp_null_allocator) // OK
+ ;
+ #pragma omp target uses_allocators(my) firstprivate(my) // { dg-error "'my' appears more than once in data clauses" }
+ ;
+ #pragma omp target private(my) uses_allocators(my) // { dg-error "'my' appears more than once in data clauses" }
+ ;
+ #pragma omp target uses_allocators(my3)
+ ;
+ #pragma omp target uses_allocators(s.a1)
+ // { dg-error "expected '\\)' before '\\.' token" "" { target *-*-* } .-1 }
+ // { dg-error "allocator 's' must be of 'omp_allocator_handle_t' type" "" { target *-*-* } .-2 }
+ ;
+ #pragma omp target uses_allocators(my4)
+ // { dg-error "allocator 'my4' must be of 'omp_allocator_handle_t' type" "" { target *-*-* } .-1 }
+ ;
+ #pragma omp target uses_allocators(my4[0])
+ // { dg-error "expected '\\)' before '\\\[' token" "" { target *-*-* } .-1 }
+ // { dg-error "allocator 'my4' must be of 'omp_allocator_handle_t' type" "" { target *-*-* } .-2 }
+ ;
+ #pragma omp target uses_allocators(memspace(omp_default_mem_space) : my, my(t)) // { dg-error "legacy 'my\\(t\\)' traits syntax not allowed in 'uses_allocators' clause when using modifiers" }
+ ;
+}
+
+// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 57 }
+// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 61 }
+// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 63 }
+// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 76 }
--- /dev/null
+/* { dg-additional-options "-fdiagnostics-show-caret -Wdeprecated-openmp" } */
+
+typedef enum omp_allocator_handle_t
+#if __cplusplus >= 201103L
+: __UINTPTR_TYPE__
+#endif
+{
+ __omp_allocator_handle_t_max__ = __UINTPTR_MAX__
+} omp_allocator_handle_t;
+
+typedef struct omp_alloctrait_t
+{
+ int key;
+ int value;
+} omp_alloctrait_t;
+
+
+void f()
+{
+ omp_allocator_handle_t myalloc;
+ const omp_alloctrait_t mytraits[] = {};
+ #pragma omp target uses_allocators(myalloc(mytraits))
+/* { dg-begin-multiline-output "" }
+ #pragma omp target uses_allocators(myalloc(mytraits))
+ ^~~~~~~~~~~~~~~~~
+ -------
+ traits : myalloc
+ { dg-end-multiline-output "" } */
+ ;
+
+// { dg-warning "38: the specification of arguments to 'uses_allocators' where each item is of the form 'allocator\\(traits\\)' is deprecated since OpenMP 5.2 \\\[-Wdeprecated-openmp\\\]" "" { target *-*-* } 22 }
+
+
+ #pragma omp target uses_allocators ( myalloc ( mytraits ) )
+/* { dg-begin-multiline-output "" }
+ #pragma omp target uses_allocators ( myalloc ( mytraits ) )
+ ^~~~~~~~~~~~~~~~~~~~~~~
+ -------
+ traits : myalloc
+ { dg-end-multiline-output "" } */
+ ;
+
+// { dg-warning "42: the specification of arguments to 'uses_allocators' where each item is of the form 'allocator\\(traits\\)' is deprecated since OpenMP 5.2 \\\[-Wdeprecated-openmp\\\]" "" { target *-*-* } 34 }
+}
+
+// { dg-excess-errors "sorry, unimplemented: 'uses_allocators' clause" }
+// { dg-excess-errors "sorry, unimplemented: 'uses_allocators' clause" }
--- /dev/null
+// { dg-do compile }
+/* { dg-additional-options "-Wno-deprecated-openmp" } */
+
+//#include <omp.h>
+
+typedef __UINTPTR_TYPE__ omp_uintptr_t;
+
+#if __cplusplus >= 201103L
+# define __GOMP_UINTPTR_T_ENUM : omp_uintptr_t
+#else
+# define __GOMP_UINTPTR_T_ENUM
+#endif
+
+typedef enum omp_memspace_handle_t __GOMP_UINTPTR_T_ENUM
+{
+ omp_default_mem_space = 0,
+ omp_large_cap_mem_space = 1,
+ omp_const_mem_space = 2,
+ omp_high_bw_mem_space = 3,
+ omp_low_lat_mem_space = 4,
+ ompx_gnu_managed_mem_space = 200,
+ __omp_memspace_handle_t_max__ = __UINTPTR_MAX__
+} omp_memspace_handle_t;
+
+typedef enum omp_allocator_handle_t __GOMP_UINTPTR_T_ENUM
+{
+ omp_null_allocator = 0,
+ omp_default_mem_alloc = 1,
+ omp_large_cap_mem_alloc = 2,
+ omp_const_mem_alloc = 3,
+ omp_high_bw_mem_alloc = 4,
+ omp_low_lat_mem_alloc = 5,
+ omp_cgroup_mem_alloc = 6,
+ omp_pteam_mem_alloc = 7,
+ omp_thread_mem_alloc = 8,
+ ompx_gnu_pinned_mem_alloc = 200,
+ ompx_gnu_managed_mem_alloc = 201,
+ __omp_allocator_handle_t_max__ = __UINTPTR_MAX__
+} omp_allocator_handle_t;
+
+typedef struct omp_alloctrait_t
+{
+// omp_alloctrait_key_t key;
+// omp_uintptr_t value;
+} omp_alloctrait_t;
+
+
+template<typename TH, TH alloc>
+void f()
+{
+ #pragma omp target uses_allocators(alloc)
+ ;
+}
+
+template<typename TH, typename TT>
+void g(TT trait, TH alloc)
+{
+ #pragma omp target uses_allocators(alloc)
+ ;
+ #pragma omp target uses_allocators(alloc(trait))
+ // { dg-error "traits array 'trait' must be defined in same scope as the construct on which the clause appears" "" { target *-*-* } .-1 }
+ // { dg-error "traits array 'trait' must be of 'const omp_alloctrait_t \\\[\\\]' type" "" { target *-*-* } .-2 }
+ ;
+}
+
+void use()
+{
+ omp_allocator_handle_t my;
+ static const omp_alloctrait_t t[] = {};
+
+ f<omp_allocator_handle_t, omp_null_allocator>(); // OK
+ f<omp_allocator_handle_t, omp_default_mem_alloc>(); // OK
+
+ g<omp_allocator_handle_t, const omp_alloctrait_t[]>(t, my); // 't'/traits not in the same scope
+}
+
+template<typename TH, TH alloc>
+void f2()
+{
+ #pragma omp target uses_allocators(alloc)
+ // { dg-error "allocator '\\(omp_allocator_handle_t\\)300' must be either a variable or a predefined allocator" "" { target *-*-* } .-1 }
+ ;
+}
+
+template<typename TH, typename TT>
+void g2(TH alloc)
+{
+ TT t = {};
+ #pragma omp target uses_allocators(alloc(t))
+ ;
+}
+
+void use2()
+{
+ omp_allocator_handle_t my;
+ const omp_allocator_handle_t wrong = (omp_allocator_handle_t) 300;
+
+ f2<omp_allocator_handle_t, wrong>(); // 300 is not a predefined allocator
+
+ g2<omp_allocator_handle_t, const omp_alloctrait_t[]>(my); // OK
+}
+
+// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 51 }
+// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 58 }
+// { dg-message "sorry, unimplemented: 'uses_allocators' clause" "" { target *-*-* } 89 }
--- /dev/null
+/* { dg-additional-options "-fdiagnostics-show-caret -Wdeprecated-openmp" } */
+
+typedef enum omp_allocator_handle_t
+#if __cplusplus >= 201103L
+: __UINTPTR_TYPE__
+#endif
+{
+ __omp_allocator_handle_t_max__ = __UINTPTR_MAX__
+} omp_allocator_handle_t;
+
+typedef struct omp_alloctrait_t
+{
+ int key;
+ int value;
+} omp_alloctrait_t;
+
+
+void f()
+{
+ omp_allocator_handle_t myalloc;
+ const omp_alloctrait_t mytraits[] = {};
+ #pragma omp target uses_allocators(myalloc(mytraits))
+/* { dg-begin-multiline-output "" }
+ #pragma omp target uses_allocators(myalloc(mytraits))
+ ^~~~~~~~~~~~~~~~~
+ -------
+ traits : myalloc
+ { dg-end-multiline-output "" } */
+ ;
+
+// { dg-warning "38: the specification of arguments to 'uses_allocators' where each item is of the form 'allocator\\(traits\\)' is deprecated since OpenMP 5.2 \\\[-Wdeprecated-openmp\\\]" "" { target *-*-* } 22 }
+
+
+ #pragma omp target uses_allocators ( myalloc ( mytraits ) )
+/* { dg-begin-multiline-output "" }
+ #pragma omp target uses_allocators ( myalloc ( mytraits ) )
+ ^~~~~~~~~~~~~~~~~~~~~~~
+ -------
+ traits : myalloc
+ { dg-end-multiline-output "" } */
+ ;
+
+// { dg-warning "42: the specification of arguments to 'uses_allocators' where each item is of the form 'allocator\\(traits\\)' is deprecated since OpenMP 5.2 \\\[-Wdeprecated-openmp\\\]" "" { target *-*-* } 34 }
+}
+
+// { dg-excess-errors "sorry, unimplemented: 'uses_allocators' clause" }
+// { dg-excess-errors "sorry, unimplemented: 'uses_allocators' clause" }
/* OpenMP clause: dyn_groupprivate ( [fallback (...)] : integer-expression). */
OMP_CLAUSE_DYN_GROUPPRIVATE,
+
+ /* OpenMP clause: uses_allocators. */
+ OMP_CLAUSE_USES_ALLOCATORS,
};
#undef DEFTREESTRUCT
pp_right_paren (pp);
break;
+ case OMP_CLAUSE_USES_ALLOCATORS:
+ pp_string (pp, "uses_allocators(");
+ dump_generic_node (pp, OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (clause),
+ spc, flags, false);
+ pp_string (pp, ": memspace(");
+ dump_generic_node (pp, OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (clause),
+ spc, flags, false);
+ pp_string (pp, "), traits(");
+ dump_generic_node (pp, OMP_CLAUSE_USES_ALLOCATORS_TRAITS (clause),
+ spc, flags, false);
+ pp_right_paren (pp);
+ pp_right_paren (pp);
+ break;
+
case OMP_CLAUSE_AFFINITY:
pp_string (pp, "affinity(");
{
1, /* OMP_CLAUSE_NOVARIANTS */
1, /* OMP_CLAUSE_NOCONTEXT */
1, /* OMP_CLAUSE_DYN_GROUPPRIVATE */
+ 3, /* OMP_CLAUSE_USES_ALLOCATORS */
};
const char * const omp_clause_code_name[] =
"novariants",
"nocontext",
"dyn_groupprivate",
+ "uses_allocators",
};
/* Unless specific to OpenACC, we tend to internally maintain OpenMP-centric
#define OMP_CLAUSE_ALLOCATE_COMBINED(NODE) \
(OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_ALLOCATE)->base.public_flag)
+#define OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR(NODE) \
+ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_USES_ALLOCATORS), 0)
+
+#define OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE(NODE) \
+ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_USES_ALLOCATORS), 1)
+
+#define OMP_CLAUSE_USES_ALLOCATORS_TRAITS(NODE) \
+ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_USES_ALLOCATORS), 2)
+
#define OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR(NODE) \
OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_NUM_TEAMS), 0)
--- /dev/null
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program main
+ use iso_c_binding
+ use omp_lib
+ implicit none (type, external)
+ integer :: x, xbuf(10)
+ integer(c_intptr_t) :: iptr
+ integer(omp_allocator_handle_kind) :: my_alloc
+ type(omp_alloctrait), parameter :: trait(*) = [omp_alloctrait(omp_atk_alignment, 128)]
+
+ !$omp target uses_allocators(omp_low_lat_mem_alloc) map(tofrom: x, xbuf) defaultmap(none)
+ !$omp parallel allocate(allocator(omp_low_lat_mem_alloc), align(128): x, xbuf) if(.false.) firstprivate(x, xbuf)
+ if (mod (TRANSFER (loc(x), iptr), 128) /= 0) &
+ stop 1
+ if (mod (TRANSFER (loc(xbuf), iptr), 128) /= 0) &
+ stop 2
+ !$omp end parallel
+ !$omp end target
+
+ my_alloc = transfer(int(z'ABCD', omp_allocator_handle_kind), my_alloc)
+
+ !$omp target uses_allocators(traits(trait): my_alloc) defaultmap(none) map(tofrom: x, xbuf)
+ !$omp parallel allocate(allocator(my_alloc): x, xbuf) if(.false.) firstprivate(x, xbuf)
+ if (mod (TRANSFER (loc(x), iptr), 128) /= 0) &
+ stop 3
+ if (mod (TRANSFER (loc(xbuf), iptr), 128) /= 0) &
+ stop 4
+ !$omp end parallel
+ !$omp end target
+
+ if (transfer(my_alloc, 0_omp_allocator_handle_kind) /= int(z'ABCD', omp_allocator_handle_kind)) &
+ stop 5
+
+ ! The following creates an allocator with empty traits + default mem space.
+ !$omp target uses_allocators(my_alloc) map(tofrom: x, xbuf) defaultmap(none)
+ !$omp parallel allocate(allocator(my_alloc), align(128): x, xbuf) if(.false.) firstprivate(x, xbuf)
+ if (mod (TRANSFER (loc(x), iptr), 128) /= 0) &
+ stop 6
+ if (mod (TRANSFER (loc(xbuf), iptr), 128) /= 0) &
+ stop 7
+ !$omp end parallel
+ !$omp end target
+
+ if (transfer(my_alloc, 0_omp_allocator_handle_kind) /= int(z'ABCD', omp_allocator_handle_kind)) &
+ stop 8
+end
+
+
+! FIXME ENABLE: 'dg FIXME final' -> 'dg-final'
+! { dg FIXME final { scan-tree-dump-times "#pragma omp target .*private\\(my_alloc\\).*uses_allocators\\(my_alloc: memspace\\(\\), traits\\(trait\\)\\)" 1 "gimple" } }
+! { dg FIXME final { scan-tree-dump-times "#pragma omp target .*private\\(my_alloc\\).*uses_allocators\\(my_alloc: memspace\\(\\), traits\\(\\)\\)" 1 "gimple" } }
+
+! FIXME ENABLE code above for "gimple" once it has been implemented:
+! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" "" { target *-*-* } 23 }
+! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" "" { target *-*-* } 36 }
+! { dg-bogus "'my_alloc' not specified in enclosing 'target'" "bogus issue because clause is ignored" { xfail *-*-* } 24 }
+! { dg-bogus "'my_alloc' not specified in enclosing 'target'" "bogus issue because clause is ignored" { xfail *-*-* } 37 }
!$omp target uses_allocators(a1 (trait3)) ! { dg-error "Traits array 'trait3' in USES_ALLOCATORS .1. must be a one-dimensional named constant array of type 'omp_alloctrait'" }
block; end block
end
+
+subroutine null_allocator_ok
+ use omp_lib
+ implicit none
+ integer(omp_allocator_handle_kind) :: my, my2
+ integer(omp_allocator_handle_kind), parameter :: my3 = -9
+ !$omp target uses_allocators(my, my2) ! OK -> default settings
+ block; end block
+ !$omp target uses_allocators(my3) ! { dg-error "'my3' at .1. in USES_ALLOCATORS must either a variable or a predefined allocator" }
+ block; end block
+ !$omp target uses_allocators(omp_default_mem_alloc, omp_null_allocator) ! OK -> omp_null_allocator
+ block; end block
+ !$omp target uses_allocators(my, omp_null_allocator) firstprivate(my) ! { dg-error "Symbol 'my' present on both data and map clauses" }
+ block; end block
+end
projects / thirdparty / gcc.git / commitdiff
