@@ -283,6 +283,7 @@ DEF_FUNCTION_TYPE_1 (BT_FN_DFLOAT32_DFLOAT32, BT_DFLOAT32, BT_DFLOAT32)
DEF_FUNCTION_TYPE_1 (BT_FN_DFLOAT64_DFLOAT64, BT_DFLOAT64, BT_DFLOAT64)
DEF_FUNCTION_TYPE_1 (BT_FN_DFLOAT128_DFLOAT128, BT_DFLOAT128, BT_DFLOAT128)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_VPTR, BT_VOID, BT_VOLATILE_PTR)
+DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRMODE, BT_VOID, BT_PTRMODE)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRPTR, BT_VOID, BT_PTR_PTR)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_CONST_PTR, BT_VOID, BT_CONST_PTR)
DEF_FUNCTION_TYPE_1 (BT_FN_UINT_UINT, BT_UINT, BT_UINT)
@@ -641,6 +642,8 @@ DEF_FUNCTION_TYPE_3 (BT_FN_PTR_SIZE_SIZE_PTRMODE,
BT_PTR, BT_SIZE, BT_SIZE, BT_PTRMODE)
DEF_FUNCTION_TYPE_3 (BT_FN_VOID_PTR_UINT8_PTRMODE, BT_VOID, BT_PTR, BT_UINT8,
BT_PTRMODE)
+DEF_FUNCTION_TYPE_3 (BT_FN_PTRMODE_PTRMODE_INT_PTR, BT_PTRMODE, BT_PTRMODE,
+ BT_INT, BT_PTR)
DEF_FUNCTION_TYPE_4 (BT_FN_SIZE_CONST_PTR_SIZE_SIZE_FILEPTR,
BT_SIZE, BT_CONST_PTR, BT_SIZE, BT_SIZE, BT_FILEPTR)
@@ -1870,6 +1870,7 @@ c_omp_split_clauses (location_t loc, enum tree_code code,
case OMP_CLAUSE_HAS_DEVICE_ADDR:
case OMP_CLAUSE_DEFAULTMAP:
case OMP_CLAUSE_DEPEND:
+ case OMP_CLAUSE_USES_ALLOCATORS:
s = C_OMP_CLAUSE_SPLIT_TARGET;
break;
case OMP_CLAUSE_NUM_TEAMS:
@@ -154,6 +154,7 @@ enum pragma_omp_clause {
PRAGMA_OMP_CLAUSE_UNTIED,
PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR,
PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR,
+ PRAGMA_OMP_CLAUSE_USES_ALLOCATORS,
/* Clauses for OpenACC. */
PRAGMA_OACC_CLAUSE_ASYNC,
@@ -12907,6 +12907,8 @@ c_parser_omp_clause_name (c_parser *parser)
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))
@@ -15624,6 +15626,225 @@ c_parser_omp_clause_allocate (c_parser *parser, tree list)
return nl;
}
+/* OpenMP 5.2:
+ uses_allocators ( allocator-list )
+
+ allocator-list:
+ allocator
+ allocator , allocator-list
+ allocator ( traits-array )
+ allocator ( traits-array ) , allocator-list
+
+ uses_allocators ( modifier : allocator )
+ uses_allocators ( modifier , modifier : allocator )
+
+ 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;
+ matching_parens parens;
+ if (!parens.require_open (parser))
+ return list;
+
+ bool has_modifiers = false;
+ tree memspace_expr = NULL_TREE;
+ tree traits_var = NULL_TREE;
+
+ if (c_parser_next_token_is (parser, CPP_NAME))
+ {
+ c_token *tok = c_parser_peek_token (parser);
+ const char *p = IDENTIFIER_POINTER (tok->value);
+
+ if (strcmp ("traits", p) == 0 || strcmp ("memspace", p) == 0)
+ {
+ has_modifiers = true;
+ c_parser_consume_token (parser);
+ matching_parens parens2;;
+ parens2.require_open (parser);
+
+ if (c_parser_next_token_is (parser, CPP_NAME)
+ && (c_parser_peek_token (parser)->id_kind == C_ID_ID
+ || c_parser_peek_token (parser)->id_kind == C_ID_TYPENAME))
+ {
+ tok = c_parser_peek_token (parser);
+ tree t = lookup_name (tok->value);
+
+ if (t == NULL_TREE)
+ {
+ undeclared_variable (tok->location, tok->value);
+ t = error_mark_node;
+ }
+ else
+ {
+ if (strcmp ("memspace", p) == 0)
+ memspace_expr = t;
+ else
+ traits_var = t;
+ }
+ c_parser_consume_token (parser);
+ }
+
+ parens2.skip_until_found_close (parser);
+
+ if (c_parser_next_token_is (parser, CPP_COMMA))
+ {
+ c_parser_consume_token (parser);
+ 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, "memspace") != 0 && strcmp (q, "traits") != 0)
+ {
+ c_parser_error (parser, "expected %<memspace%> or %<traits%>");
+ 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;
+ }
+
+ if (c_parser_next_token_is (parser, CPP_NAME)
+ && (c_parser_peek_token (parser)->id_kind == C_ID_ID
+ || c_parser_peek_token (parser)->id_kind == C_ID_TYPENAME))
+ {
+ tok = c_parser_peek_token (parser);
+ tree t = lookup_name (tok->value);
+ if (t == NULL_TREE)
+ {
+ undeclared_variable (tok->location, tok->value);
+ t = error_mark_node;
+ }
+ else
+ {
+ if (strcmp ("memspace", q) == 0)
+ memspace_expr = t;
+ else
+ traits_var = t;
+ }
+ c_parser_consume_token (parser);
+ }
+ parens2.skip_until_found_close (parser);
+ }
+ }
+ }
+
+ if (has_modifiers)
+ {
+ if (!c_parser_require (parser, CPP_COLON, "expected %<:%>"))
+ {
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+
+ if (c_parser_next_token_is (parser, CPP_NAME)
+ && c_parser_peek_token (parser)->id_kind == C_ID_ID)
+ {
+ tree t = lookup_name (c_parser_peek_token (parser)->value);
+
+ if (t == NULL_TREE)
+ {
+ undeclared_variable (c_parser_peek_token (parser)->location,
+ c_parser_peek_token (parser)->value);
+ t = error_mark_node;
+ }
+ else if (t != error_mark_node)
+ {
+ 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) = traits_var;
+ OMP_CLAUSE_CHAIN (c) = list;
+
+ nl = c;
+ }
+ c_parser_consume_token (parser);
+
+ if (c_parser_next_token_is (parser, CPP_OPEN_PAREN))
+ c_parser_error (parser, "modifiers cannot be used with "
+ "legacy array syntax");
+ }
+ else
+ c_parser_error (parser, "expected identifier");
+ }
+ else
+ {
+ traits_var = NULL_TREE;
+
+ nl = list;
+ while (true)
+ {
+ if (c_parser_next_token_is (parser, CPP_NAME)
+ && c_parser_peek_token (parser)->id_kind == C_ID_ID)
+ {
+ c_token *tok = c_parser_peek_token (parser);
+ tree t = lookup_name (tok->value);
+
+ if (t == NULL_TREE)
+ {
+ undeclared_variable (tok->location, tok->value);
+ break;
+ }
+ else if (t == error_mark_node)
+ break;
+
+ c_parser_consume_token (parser);
+
+ if (c_parser_next_token_is (parser, CPP_OPEN_PAREN))
+ {
+ matching_parens parens2;
+ parens2.consume_open (parser);
+ if (c_parser_next_token_is (parser, CPP_NAME)
+ && c_parser_peek_token (parser)->id_kind == C_ID_ID)
+ {
+ tok = c_parser_peek_token (parser);
+ traits_var = lookup_name (tok->value);
+ if (traits_var == NULL_TREE)
+ {
+ undeclared_variable (tok->location, tok->value);
+ break;
+ }
+ c_parser_consume_token (parser);
+ }
+ else
+ c_parser_error (parser, "expected identifier");
+ parens2.require_close (parser);
+ }
+
+ tree c = build_omp_clause (clause_loc,
+ OMP_CLAUSE_USES_ALLOCATORS);
+ OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c) = t;
+ OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c) = NULL_TREE;
+ OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c) = traits_var;
+ OMP_CLAUSE_CHAIN (c) = nl;
+
+ nl = c;
+ }
+
+ if (c_parser_next_token_is_not (parser, CPP_COMMA))
+ break;
+ c_parser_consume_token (parser);
+ }
+ }
+
+ parens.skip_until_found_close (parser);
+ return nl;
+}
+
/* OpenMP 4.0:
linear ( variable-list )
linear ( variable-list : expression )
@@ -17050,6 +17271,10 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_omp_clause_allocate (parser, clauses);
c_name = "allocate";
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_LINEAR:
clauses = c_parser_omp_clause_linear (parser, clauses);
c_name = "linear";
@@ -21061,7 +21286,8 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_THREAD_LIMIT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)\
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR))
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR)\
+ | (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)
@@ -14763,6 +14763,102 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
break;
+ case OMP_CLAUSE_USES_ALLOCATORS:
+ t = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);
+ if (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;
+ }
+ 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 must be of %<omp_allocator_handle_t%> type");
+ remove = true;
+ }
+ if (TREE_CODE (t) == CONST_DECL)
+ {
+ if (OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c)
+ || OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c))
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "modifiers cannot be used with pre-defined "
+ "allocators");
+
+ /* Currently for pre-defined allocators in libgomp, we do not
+ require additional init/fini inside target regions, so discard
+ such clauses. */
+ remove = true;
+ }
+ t = OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c);
+ if (t != NULL_TREE
+ && (TREE_CODE (t) != CONST_DECL
+ || 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 must be "
+ "constant enum of %<omp_memspace_handle_t%> type");
+ remove = true;
+ }
+ t = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c);
+ if (t != NULL_TREE)
+ {
+ bool type_err = false;
+
+ if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE)
+ 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 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 "
+ "of constant values");
+
+ remove = true;
+ }
+ }
+ }
+
+ if (remove)
+ break;
+ else
+ {
+ /* Create a private clause for the allocator variable, placed
+ prior to current uses_allocators clause. */
+ tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_PRIVATE);
+ OMP_CLAUSE_DECL (nc) = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);
+ OMP_CLAUSE_CHAIN (nc) = c;
+ *pc = nc;
+
+ pc = &OMP_CLAUSE_CHAIN (c);
+ continue;
+ }
+
case OMP_CLAUSE_DEPEND:
t = OMP_CLAUSE_DECL (c);
if (t == NULL_TREE)
@@ -36490,6 +36490,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
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))
@@ -38733,6 +38735,246 @@ cp_parser_omp_clause_allocate (cp_parser *parser, tree list)
return nlist;
}
+/* OpenMP 5.2:
+ uses_allocators ( allocator-list )
+
+ allocator-list:
+ allocator
+ allocator , allocator-list
+ allocator ( traits-array )
+ allocator ( traits-array ) , allocator-list
+
+ uses_allocators ( modifier : allocator )
+ uses_allocators ( modifier , modifier : allocator )
+
+ 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 t = NULL_TREE, nl;
+ matching_parens parens;
+ if (!parens.require_open (parser))
+ return list;
+
+ bool has_modifiers = false;
+ tree memspace_expr = NULL_TREE;
+ tree traits_var = NULL_TREE;
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+ {
+ cp_token *tok = cp_lexer_peek_token (parser->lexer);
+ const char *p = IDENTIFIER_POINTER (tok->u.value);
+
+ if (strcmp ("traits", p) == 0 || strcmp ("memspace", p) == 0)
+ {
+ has_modifiers = true;
+ cp_lexer_consume_token (parser->lexer);
+ matching_parens parens2;;
+ parens2.require_open (parser);
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+ {
+ tok = cp_lexer_peek_token (parser->lexer);
+ tree id = tok->u.value;
+
+ t = cp_parser_lookup_name_simple (parser, id, tok->location);
+ if (t == error_mark_node)
+ cp_parser_name_lookup_error (parser, id, t, NLE_NULL,
+ tok->location);
+ else
+ {
+ if (strcmp ("memspace", p) == 0)
+ memspace_expr = t;
+ else
+ traits_var = t;
+ }
+ cp_lexer_consume_token (parser->lexer);
+ }
+
+ if (t == error_mark_node || !parens.require_close (parser))
+ {
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return list;
+ }
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
+ {
+ cp_lexer_consume_token (parser->lexer);
+ tok = cp_lexer_peek_token (parser->lexer);
+ const char *q = "";
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+ q = IDENTIFIER_POINTER (tok->u.value);
+
+ if (strcmp (q, "memspace") != 0 && strcmp (q, "traits") != 0)
+ {
+ cp_parser_error (parser, "expected %<memspace%> or %<traits%>");
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return list;
+ }
+ else if (strcmp (p, q) == 0)
+ {
+ error_at (tok->location, "duplicate %qs modifier", p);
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return list;
+ }
+ cp_lexer_consume_token (parser->lexer);
+ if (!parens2.require_open (parser))
+ {
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return list;
+ }
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+ {
+ tok = cp_lexer_peek_token (parser->lexer);
+ tree id = tok->u.value;
+
+ t = cp_parser_lookup_name_simple (parser, id, tok->location);
+ if (t == error_mark_node)
+ cp_parser_name_lookup_error (parser, id, t, NLE_NULL,
+ tok->location);
+ else
+ {
+ if (strcmp ("memspace", q) == 0)
+ memspace_expr = t;
+ else
+ traits_var = t;
+ }
+ cp_lexer_consume_token (parser->lexer);
+ }
+
+ if (t == error_mark_node || !parens.require_close (parser))
+ {
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return list;
+ }
+ }
+
+ }
+ }
+
+ if (has_modifiers)
+ {
+ if (!cp_parser_require (parser, CPP_COLON, RT_COLON))
+ {
+ cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return list;
+ }
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+ {
+ cp_token *tok = cp_lexer_peek_token (parser->lexer);
+ tree id = tok->u.value;
+ tree t = cp_parser_lookup_name_simple (parser, id, tok->location);
+
+ if (t == error_mark_node)
+ cp_parser_name_lookup_error (parser, id, t, NLE_NULL,
+ tok->location);
+ else
+ {
+ 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) = traits_var;
+ OMP_CLAUSE_CHAIN (c) = list;
+
+ nl = c;
+ }
+ cp_lexer_consume_token (parser->lexer);
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_PAREN))
+ cp_parser_error (parser, "modifiers cannot be used with "
+ "legacy array syntax");
+ }
+ else
+ cp_parser_error (parser, "expected identifier");
+ }
+ else
+ {
+ traits_var = NULL_TREE;
+
+ nl = list;
+ while (true)
+ {
+ if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+ {
+ cp_token *tok = cp_lexer_peek_token (parser->lexer);
+ tree id = tok->u.value;
+
+ tree t = cp_parser_lookup_name_simple (parser, id, tok->location);
+
+ if (t == error_mark_node)
+ cp_parser_name_lookup_error (parser, id, t, NLE_NULL,
+ tok->location);
+ cp_lexer_consume_token (parser->lexer);
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_PAREN))
+ {
+ matching_parens parens2;
+ parens2.consume_open (parser);
+ if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+ {
+ tok = cp_lexer_peek_token (parser->lexer);
+ id = tok->u.value;
+ traits_var = cp_parser_lookup_name_simple (parser, id,
+ tok->location);
+ if (traits_var == error_mark_node)
+ {
+ cp_parser_name_lookup_error (parser, id, traits_var,
+ NLE_NULL, tok->location);
+ break;
+ }
+ cp_lexer_consume_token (parser->lexer);
+ }
+ else
+ cp_parser_error (parser, "expected identifier");
+ parens2.require_close (parser);
+ }
+
+ tree c = build_omp_clause (clause_loc,
+ OMP_CLAUSE_USES_ALLOCATORS);
+ OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c) = t;
+ OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c) = NULL_TREE;
+ OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c) = traits_var;
+ OMP_CLAUSE_CHAIN (c) = nl;
+
+ nl = c;
+ }
+
+ if (cp_lexer_next_token_is_not (parser->lexer, CPP_COMMA))
+ break;
+ cp_lexer_consume_token (parser->lexer);
+ }
+ }
+
+ parens.require_close (parser);
+ return nl;
+}
+
/* OpenMP 2.5:
lastprivate ( variable-list )
@@ -40283,6 +40525,10 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
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;
@@ -44291,7 +44537,8 @@ cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_THREAD_LIMIT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)\
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR))
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR)\
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USES_ALLOCATORS))
static bool
cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
@@ -7772,6 +7772,90 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
goto handle_field_decl;
+ case OMP_CLAUSE_USES_ALLOCATORS:
+ t = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);
+ 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 must be of %<omp_allocator_handle_t%> type");
+ remove = true;
+ }
+ if (TREE_CODE (t) == CONST_DECL)
+ {
+ if (OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c)
+ || OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c))
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "modifiers cannot be used with pre-defined "
+ "allocators");
+
+ /* Currently for pre-defined allocators in libgomp, we do not
+ require additional init/fini inside target regions, so discard
+ such clauses. */
+ remove = true;
+ }
+ t = OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c);
+ if (t != NULL_TREE
+ && (TREE_CODE (t) != CONST_DECL
+ || 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 must be "
+ "constant enum of %<omp_memspace_handle_t%> type");
+ remove = true;
+ }
+ t = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c);
+ if (t != NULL_TREE)
+ {
+ bool type_err = false;
+
+ if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE)
+ 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 must be of "
+ "%<const omp_alloctrait_t []%> type");
+ remove = true;
+ }
+ else
+ {
+ tree cst_val = decl_constant_value (t);
+ if (cst_val == t)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "traits array must be "
+ "of constant values");
+
+ remove = true;
+ }
+ }
+ }
+ if (remove)
+ break;
+ else
+ {
+ /* Create a private clause for the allocator variable, placed
+ prior to current uses_allocators clause. */
+ tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_PRIVATE);
+ OMP_CLAUSE_DECL (nc) = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);
+ OMP_CLAUSE_CHAIN (nc) = c;
+ *pc = nc;
+
+ pc = &OMP_CLAUSE_CHAIN (c);
+ continue;
+ }
+
case OMP_CLAUSE_DEPEND:
t = OMP_CLAUSE_DECL (c);
if (t == NULL_TREE)
@@ -1353,6 +1353,8 @@ typedef struct gfc_omp_namelist
struct gfc_omp_namelist_udr *udr;
gfc_namespace *ns;
} u2;
+ struct gfc_symbol *memspace_sym;
+ struct gfc_symbol *traits_sym;
struct gfc_omp_namelist *next;
locus where;
}
@@ -1394,6 +1396,7 @@ enum
OMP_LIST_NONTEMPORAL,
OMP_LIST_ALLOCATE,
OMP_LIST_HAS_DEVICE_ADDR,
+ OMP_LIST_USES_ALLOCATORS,
OMP_LIST_NUM /* Must be the last. */
};
@@ -948,6 +948,7 @@ enum omp_mask2
OMP_CLAUSE_ATTACH,
OMP_CLAUSE_NOHOST,
OMP_CLAUSE_HAS_DEVICE_ADDR, /* OpenMP 5.1 */
+ OMP_CLAUSE_USES_ALLOCATORS, /* OpenMP 5.2 */
/* This must come last. */
OMP_MASK2_LAST
};
@@ -1364,6 +1365,238 @@ gfc_match_omp_clause_reduction (char pc, gfc_omp_clauses *c, bool openacc,
return MATCH_YES;
}
+/* uses_allocators ( allocator-list )
+
+ allocator-list:
+ allocator
+ allocator , allocator-list
+ allocator ( traits-array )
+ allocator ( traits-array ) , allocator-list
+
+ uses_allocators ( modifier : allocator )
+ uses_allocators ( modifier , modifier : allocator )
+
+ modifier:
+ traits ( traits-array )
+ memspace ( mem-space-handle ) */
+
+static match
+gfc_match_omp_clause_uses_allocators (gfc_omp_clauses *c)
+{
+ gfc_symbol *sym;
+ gfc_symbol *memspace_sym= NULL;
+ gfc_symbol *traits_sym= NULL;
+ bool memspace_seen = false, traits_seen = false;
+ match m;
+ int i = 0;
+
+ if (gfc_match ("uses_allocators ( ") != MATCH_YES)
+ return MATCH_NO;
+
+ gfc_symbol *allocator_handle_kind, * memspace_handle_kind;
+
+ if (gfc_get_symbol ("omp_allocator_handle_kind", NULL, &sym)
+ || !sym->value
+ || sym->value->expr_type != EXPR_CONSTANT
+ || sym->value->ts.type != BT_INTEGER)
+ {
+ gfc_error ("OpenMP %<omp_allocator_handle_kind%> constant not found by "
+ "%<uses_allocators%> clause at %C");
+ goto error;
+ }
+ allocator_handle_kind = sym;
+
+ if (gfc_get_symbol ("omp_memspace_handle_kind", NULL, &sym)
+ || !sym->value
+ || sym->value->expr_type != EXPR_CONSTANT
+ || sym->value->ts.type != BT_INTEGER)
+ {
+ gfc_error ("OpenMP %<omp_memspace_handle_kind%> constant not found by "
+ "%<uses_allocators%> clause at %C");
+ goto error;
+ }
+ memspace_handle_kind = sym;
+
+ do
+ {
+ if (++i > 2)
+ {
+ gfc_error ("Only two modifiers are allowed on %<uses_allocators%> "
+ "clause at %C");
+ goto error;
+ }
+
+ if (gfc_match ("memspace ( ") == MATCH_YES)
+ {
+ if (memspace_seen)
+ {
+ gfc_error ("Multiple memspace modifiers at %C");
+ goto error;
+ }
+ memspace_seen = true;
+ m = gfc_match_symbol (&sym, 1);
+ if (m == MATCH_YES)
+ memspace_sym = sym;
+ else
+ goto error;
+ if (gfc_match_char (')') != MATCH_YES)
+ goto error;
+ }
+ else if (gfc_match ("traits ( ") == MATCH_YES)
+ {
+ if (traits_seen)
+ {
+ gfc_error ("Multiple traits modifiers at %C");
+ goto error;
+ }
+ traits_seen = true;
+ m = gfc_match_symbol (&sym, 1);
+ if (m == MATCH_YES)
+ traits_sym = sym;
+ else
+ goto error;
+ if (gfc_match_char (')') != MATCH_YES)
+ goto error;
+ }
+ else
+ break;
+ }
+ while (gfc_match (" , ") == MATCH_YES);
+
+ if ((memspace_seen || traits_seen)
+ && gfc_match (" : ") != MATCH_YES)
+ goto error;
+
+ while (true)
+ {
+ m = gfc_match_symbol (&sym, 1);
+ if (m != MATCH_YES)
+ {
+ gfc_error ("Expected name of allocator at %C");
+ goto error;
+ }
+ gfc_symbol *allocator_sym = sym;
+
+ if (gfc_match_char ('(') == MATCH_YES)
+ {
+ if (memspace_seen || traits_seen)
+ {
+ gfc_error ("Modifiers cannot be used with legacy "
+ "array syntax at %C");
+ goto error;
+ }
+ m = gfc_match_symbol (&sym, 1);
+ if (m == MATCH_YES)
+ traits_sym = sym;
+ else
+ goto error;
+ if (gfc_match_char (')') != MATCH_YES)
+ goto error;
+ }
+
+ if (traits_sym)
+ {
+ if (traits_sym->ts.type != BT_DERIVED
+ || strcmp (traits_sym->ts.u.derived->name,
+ "omp_alloctrait") != 0
+ || traits_sym->attr.flavor != FL_PARAMETER
+ || traits_sym->as->rank != 1
+ || traits_sym->value == NULL
+ || !gfc_is_constant_expr (traits_sym->value))
+ {
+ gfc_error ("%<%s%> at %C must be of constant "
+ "%<type(omp_alloctrait)%> array type and have a "
+ "constant initializer", traits_sym->name);
+ goto error;
+ }
+ gfc_set_sym_referenced (traits_sym);
+ }
+
+ if (memspace_sym)
+ {
+ const char *pos;
+ if (memspace_sym->ts.type != BT_INTEGER
+ || memspace_sym->attr.flavor != FL_PARAMETER
+ || mpz_cmp_si (memspace_handle_kind->value->value.integer,
+ memspace_sym->ts.kind) != 0
+ /* Check if identifier is of 'omp_..._mem_space' format. */
+ || (pos = strstr (memspace_sym->name, "omp_")) == NULL
+ || pos != memspace_sym->name
+ || (pos = strstr (memspace_sym->name, "_mem_space")) == NULL
+ || *(pos + strlen ("_mem_space")) != '\0')
+ {
+ gfc_error ("%<%s%> at %C is not a pre-defined memory space name",
+ memspace_sym->name);
+ goto error;
+ }
+ }
+
+ if (allocator_sym->ts.type != BT_INTEGER
+ || mpz_cmp_si (allocator_handle_kind->value->value.integer,
+ allocator_sym->ts.kind) != 0)
+ {
+ gfc_error ("%<%s%> at %C must be integer of %<%s%> kind",
+ allocator_sym->name, allocator_handle_kind->name);
+ goto error;
+ }
+
+ if (allocator_sym->attr.flavor == FL_PARAMETER)
+ {
+ const char *pos;
+ /* Check if identifier is a 'omp_..._mem_alloc' pre-defined
+ allocator. */
+ if ((pos = strstr (allocator_sym->name, "omp_")) == NULL
+ || pos != allocator_sym->name
+ || (pos = strstr (allocator_sym->name, "_mem_alloc")) == NULL
+ || *(pos + strlen ("_mem_alloc")) != '\0')
+ {
+ gfc_error ("%<%s%> at %C is not a pre-defined memory allocator",
+ allocator_sym->name);
+ goto error;
+ }
+
+ /* Currently for pre-defined allocators in libgomp, we do not
+ require additional init/fini inside target regions,
+ so do nothing here to discard such clauses. */
+ }
+ else
+ {
+ gfc_set_sym_referenced (allocator_sym);
+
+ gfc_omp_namelist *n = gfc_get_omp_namelist ();
+ n->sym = allocator_sym;
+ n->memspace_sym = memspace_sym;
+ n->traits_sym = traits_sym;
+ n->where = gfc_current_locus;
+
+ n->next = c->lists[OMP_LIST_USES_ALLOCATORS];
+ c->lists[OMP_LIST_USES_ALLOCATORS] = n;
+ }
+
+ if (gfc_match (" , ") == MATCH_YES)
+ {
+ if (memspace_seen || traits_seen)
+ {
+ gfc_error ("When using modifiers, only a single allocator can be "
+ "specified in each %<uses_allocators%> clause at %C");
+ goto error;
+ }
+ }
+ else
+ break;
+
+ memspace_sym = NULL;
+ traits_sym = NULL;
+ }
+
+ if (gfc_match_char (')') != MATCH_YES)
+ goto error;
+
+ return MATCH_YES;
+
+ error:
+ return MATCH_ERROR;
+}
/* Match with duplicate check. Matches 'name'. If expr != NULL, it
then matches '(expr)', otherwise, if open_parens is true,
@@ -2924,6 +3157,9 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
("use_device_addr (", &c->lists[OMP_LIST_USE_DEVICE_ADDR],
false, NULL, NULL, true) == MATCH_YES)
continue;
+ if ((mask & OMP_CLAUSE_USES_ALLOCATORS)
+ && gfc_match_omp_clause_uses_allocators (c) == MATCH_YES)
+ continue;
break;
case 'v':
/* VECTOR_LENGTH must be matched before VECTOR, because the latter
@@ -3650,7 +3886,7 @@ cleanup:
| OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULTMAP \
| OMP_CLAUSE_IS_DEVICE_PTR | OMP_CLAUSE_IN_REDUCTION \
| OMP_CLAUSE_THREAD_LIMIT | OMP_CLAUSE_ALLOCATE \
- | OMP_CLAUSE_HAS_DEVICE_ADDR)
+ | OMP_CLAUSE_HAS_DEVICE_ADDR | OMP_CLAUSE_USES_ALLOCATORS)
#define OMP_TARGET_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_DEVICE) | OMP_CLAUSE_MAP | OMP_CLAUSE_IF \
| OMP_CLAUSE_USE_DEVICE_PTR | OMP_CLAUSE_USE_DEVICE_ADDR)
@@ -6282,7 +6518,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
"IN_REDUCTION", "TASK_REDUCTION",
"DEVICE_RESIDENT", "LINK", "USE_DEVICE",
"CACHE", "IS_DEVICE_PTR", "USE_DEVICE_PTR", "USE_DEVICE_ADDR",
- "NONTEMPORAL", "ALLOCATE", "HAS_DEVICE_ADDR" };
+ "NONTEMPORAL", "ALLOCATE", "HAS_DEVICE_ADDR", "USES_ALLOCATORS" };
STATIC_ASSERT (ARRAY_SIZE (clause_names) == OMP_LIST_NUM);
if (omp_clauses == NULL)
@@ -6343,10 +6343,8 @@ gfc_conv_array_initializer (tree type, gfc_expr * expr)
&expr->where, flag_max_array_constructor);
return NULL_TREE;
}
- if (mpz_cmp_si (c->offset, 0) != 0)
- index = gfc_conv_mpz_to_tree (c->offset, gfc_index_integer_kind);
- else
- index = NULL_TREE;
+
+ index = gfc_conv_mpz_to_tree (c->offset, gfc_index_integer_kind);
if (mpz_cmp_si (c->repeat, 1) > 0)
{
@@ -2686,9 +2686,16 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
if (n->expr)
{
tree allocator_;
- gfc_init_se (&se, NULL);
- gfc_conv_expr (&se, n->expr);
- allocator_ = gfc_evaluate_now (se.expr, block);
+ if (n->expr->expr_type == EXPR_VARIABLE)
+ allocator_
+ = gfc_trans_omp_variable (n->expr->symtree->n.sym,
+ false);
+ else
+ {
+ gfc_init_se (&se, NULL);
+ gfc_conv_expr (&se, n->expr);
+ allocator_ = gfc_evaluate_now (se.expr, block);
+ }
OMP_CLAUSE_ALLOCATE_ALLOCATOR (node) = allocator_;
}
omp_clauses = gfc_trans_add_clause (node, omp_clauses);
@@ -3657,6 +3664,29 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
omp_clauses = gfc_trans_add_clause (node, omp_clauses);
}
break;
+ case OMP_LIST_USES_ALLOCATORS:
+ for (; n != NULL; n = n->next)
+ {
+ tree allocator = gfc_trans_omp_variable (n->sym, false);
+ tree memspace = (n->memspace_sym
+ ? gfc_conv_constant_to_tree (n->memspace_sym->value)
+ : NULL_TREE);
+ tree traits = (n->traits_sym
+ ? gfc_trans_omp_variable (n->traits_sym, false)
+ : NULL_TREE);
+
+ tree nc = build_omp_clause (input_location, OMP_CLAUSE_PRIVATE);
+ OMP_CLAUSE_DECL (nc) = allocator;
+ omp_clauses = gfc_trans_add_clause (nc, omp_clauses);
+
+ nc = build_omp_clause (input_location,
+ OMP_CLAUSE_USES_ALLOCATORS);
+ OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (nc) = allocator;
+ OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (nc) = memspace;
+ OMP_CLAUSE_USES_ALLOCATORS_TRAITS (nc) = traits;
+ omp_clauses = gfc_trans_add_clause (nc, omp_clauses);
+ }
+ break;
default:
break;
}
@@ -81,6 +81,7 @@ DEF_FUNCTION_TYPE_0 (BT_FN_UINT, BT_UINT)
DEF_FUNCTION_TYPE_0 (BT_FN_VOID, BT_VOID)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTR, BT_VOID, BT_PTR)
+DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRMODE, BT_VOID, BT_PTRMODE)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRPTR, BT_VOID, BT_PTR_PTR)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_VPTR, BT_VOID, BT_VOLATILE_PTR)
DEF_FUNCTION_TYPE_1 (BT_FN_INT_INT, BT_INT, BT_INT)
@@ -154,6 +155,8 @@ DEF_FUNCTION_TYPE_3 (BT_FN_VOID_SIZE_SIZE_PTR, BT_VOID, BT_SIZE, BT_SIZE,
DEF_FUNCTION_TYPE_3 (BT_FN_UINT_UINT_PTR_PTR, BT_UINT, BT_UINT, BT_PTR, BT_PTR)
DEF_FUNCTION_TYPE_3 (BT_FN_PTR_SIZE_SIZE_PTRMODE,
BT_PTR, BT_SIZE, BT_SIZE, BT_PTRMODE)
+DEF_FUNCTION_TYPE_3 (BT_FN_PTRMODE_PTRMODE_INT_PTR, BT_PTRMODE, BT_PTRMODE,
+ BT_INT, BT_PTR)
DEF_FUNCTION_TYPE_4 (BT_FN_VOID_OMPFN_PTR_UINT_UINT,
BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT, BT_UINT)
@@ -9148,6 +9148,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
hash_map<tree_operand_hash, tree> *struct_map_to_clause = NULL;
hash_map<tree_operand_hash, tree *> *struct_seen_clause = NULL;
hash_set<tree> *struct_deref_set = NULL;
+
+ hash_set<tree> *allocate_clauses = NULL;
+ hash_set<tree> *uses_allocators_allocators = NULL;
+
tree *prev_list_p = NULL, *orig_list_p = list_p;
int handled_depend_iterators = -1;
int nowait = -1;
@@ -9185,6 +9189,13 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
|| code == OMP_TARGET_EXIT_DATA)
omp_target_reorder_clauses (list_p);
+ if (code == OMP_TARGET
+ && (omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0)
+ {
+ allocate_clauses = new hash_set<tree> ();
+ uses_allocators_allocators = new hash_set<tree> ();
+ }
+
while ((c = *list_p) != NULL)
{
bool remove = false;
@@ -10884,6 +10895,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
OMP_CLAUSE_ALLOCATE_ALLOCATOR (c)
= get_initialized_tmp_var (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c),
pre_p, NULL, false);
+ if (allocate_clauses
+ && OMP_CLAUSE_ALLOCATE_ALLOCATOR (c)
+ && DECL_P (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c))
+ && !allocate_clauses->contains (c))
+ allocate_clauses->add (c);
+ break;
+
+ case OMP_CLAUSE_USES_ALLOCATORS:
+ decl = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);
+ if (uses_allocators_allocators
+ && !uses_allocators_allocators->contains (decl))
+ uses_allocators_allocators->add (decl);
break;
case OMP_CLAUSE_DEFAULT:
@@ -10936,6 +10959,26 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
list_p = &OMP_CLAUSE_CHAIN (c);
}
+ if (code == OMP_TARGET
+ && (omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0)
+ {
+ for (hash_set<tree>::iterator i = allocate_clauses->begin ();
+ i != allocate_clauses->end (); ++i)
+ {
+ tree allocator = OMP_CLAUSE_ALLOCATE_ALLOCATOR (*i);
+ if (uses_allocators_allocators->contains (allocator))
+ continue;
+
+ error_at (OMP_CLAUSE_LOCATION (*i),
+ "allocator %<%qE%>in %<allocate%> clause on target region "
+ "is missing %<uses_allocators(%E)%> clause",
+ DECL_NAME (allocator), DECL_NAME (allocator));
+ }
+
+ delete allocate_clauses;
+ delete uses_allocators_allocators;
+ }
+
ctx->clauses = *orig_list_p;
gimplify_omp_ctxp = ctx;
if (struct_seen_clause)
@@ -14165,6 +14208,79 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
body = NULL;
gimple_seq_add_stmt (&body, g);
}
+ else if ((ort & ORT_TARGET) != 0 && (ort & ORT_ACC) == 0)
+ {
+ gimple_seq init_seq = NULL;
+ gimple_seq fini_seq = NULL;
+
+ tree omp_init_allocator_fn
+ = builtin_decl_explicit (BUILT_IN_OMP_INIT_ALLOCATOR);
+ tree omp_destroy_allocator_fn
+ = builtin_decl_explicit (BUILT_IN_OMP_DESTROY_ALLOCATOR);
+
+ for (tree *cp = &OMP_CLAUSES (expr); *cp != NULL;)
+ if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_USES_ALLOCATORS)
+ {
+ tree c = *cp;
+ tree allocator = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);
+ tree memspace = OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c);
+ tree traits = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c);
+ tree ntraits
+ = ((traits
+ && DECL_INITIAL (traits)
+ && TREE_CODE (DECL_INITIAL (traits)) == CONSTRUCTOR)
+ ? build_int_cst (integer_type_node,
+ CONSTRUCTOR_NELTS (DECL_INITIAL (traits)))
+ : integer_zero_node);
+ tree traits_var
+ = (traits != NULL_TREE
+ ? get_initialized_tmp_var (DECL_INITIAL (traits),
+ &init_seq, NULL)
+ : null_pointer_node);
+
+ tree memspace_var = create_tmp_var (pointer_sized_int_node,
+ "memspace_enum");
+ if (memspace == NULL_TREE)
+ memspace = build_int_cst (pointer_sized_int_node, 0);
+ else
+ memspace = fold_convert (pointer_sized_int_node,
+ memspace);
+ g = gimple_build_assign (memspace_var, memspace);
+ gimple_seq_add_stmt (&init_seq, g);
+
+ /*
+ gimplify_assign (memspace_var,
+ fold_convert (pointer_sized_int_node,
+ memspace),
+ &init_seq);
+ */
+ tree initcall = build_call_expr_loc (OMP_CLAUSE_LOCATION (c),
+ omp_init_allocator_fn, 3,
+ memspace_var,
+ ntraits,
+ traits_var);
+ initcall = fold_convert (TREE_TYPE (allocator), initcall);
+ gimplify_assign (allocator, initcall, &init_seq);
+
+ g = gimple_build_call (omp_destroy_allocator_fn, 1, allocator);
+ gimple_seq_add_stmt (&fini_seq, g);
+
+ /* Finished generating runtime calls, remove USES_ALLOCATORS
+ clause. */
+ *cp = OMP_CLAUSE_CHAIN (c);
+ }
+ else
+ cp = &OMP_CLAUSE_CHAIN (*cp);
+
+ if (fini_seq)
+ {
+ gbind *bind = as_a<gbind *> (gimple_seq_first_stmt (body));
+ g = gimple_build_try (gimple_bind_body (bind),
+ fini_seq, GIMPLE_TRY_FINALLY);
+ gimple_seq_add_stmt (&init_seq, g);
+ gimple_bind_set_body (bind, init_seq);
+ }
+ }
}
else
gimplify_and_add (OMP_BODY (expr), &body);
@@ -76,6 +76,10 @@ DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_TEAM_NUM, "omp_get_team_num",
BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_NUM_TEAMS, "omp_get_num_teams",
BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_OMP_INIT_ALLOCATOR, "omp_init_allocator",
+ BT_FN_PTRMODE_PTRMODE_INT_PTR, ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_OMP_DESTROY_ALLOCATOR, "omp_destroy_allocator",
+ BT_FN_VOID_PTRMODE, ATTR_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_ATOMIC_START, "GOMP_atomic_start",
BT_FN_VOID, ATTR_NOTHROW_LEAF_LIST)
new file mode 100644
@@ -0,0 +1,46 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */
+
+#include <omp.h>
+
+omp_alloctrait_key_t k;
+omp_alloctrait_value_t v;
+
+int main (void)
+{
+ omp_allocator_handle_t foo, bar;
+ const omp_alloctrait_t foo_traits[] = { { omp_atk_pinned, omp_atv_true },
+ { omp_atk_partition, omp_atv_nearest } };
+ #pragma omp target
+ ;
+ #pragma omp target uses_allocators (bar)
+ ;
+ #pragma omp target uses_allocators (foo (foo_traits))
+ ;
+ #pragma omp target uses_allocators (foo (foo_traits), bar (foo_traits))
+ ;
+ #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo)
+ ;
+ #pragma omp target uses_allocators (traits(foo_traits) : bar)
+ ;
+ #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space), traits(foo_traits) : bar)
+ ;
+ #pragma omp target uses_allocators (traits(foo_traits), memspace(omp_high_bw_mem_space) : bar) uses_allocators(foo)
+ {
+ void *p = omp_alloc ((unsigned long) 32, bar);
+ omp_free (p, bar);
+ }
+ return 0;
+}
+
+/* { dg-final { scan-tree-dump "pragma omp target" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target private\\(bar\\) uses_allocators\\(bar: memspace\\(\\), traits\\(\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target private\\(foo\\) uses_allocators\\(foo: memspace\\(\\), traits\\(foo_traits\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target private\\(bar\\) uses_allocators\\(bar: memspace\\(\\), traits\\(foo_traits\\)\\) private\\(foo\\) uses_allocators\\(foo: memspace\\(\\), traits\\(foo_traits\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target private\\(foo\\) uses_allocators\\(foo: memspace\\(omp_high_bw_mem_space\\), traits\\(\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target private\\(bar\\) uses_allocators\\(bar: memspace\\(\\), traits\\(foo_traits\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target private\\(bar\\) uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\), traits\\(foo_traits\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target private\\(foo\\) uses_allocators\\(foo: memspace\\(\\), traits\\(\\)\\) private\\(bar\\) uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\), traits\\(foo_traits\\)\\)" "original" } } */
+
+/* { dg-final { scan-tree-dump-times "__builtin_omp_init_allocator" 9 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_omp_destroy_allocator" 9 "gimple" } } */
new file mode 100644
@@ -0,0 +1,32 @@
+/* { dg-do compile } */
+
+#include <omp.h>
+
+omp_alloctrait_key_t k;
+omp_alloctrait_value_t v;
+
+int main (void)
+{
+ omp_allocator_handle_t foo, bar;
+ const omp_alloctrait_t traits_array[] = { { omp_atk_pinned, omp_atv_true },
+ { omp_atk_partition, omp_atv_nearest } };
+ #pragma omp target uses_allocators (baz) /* { dg-error "'baz' undeclared .first use in this function." } */
+ ;
+ #pragma omp target uses_allocators (foo (xyz)) /* { dg-error "'xyz' undeclared .first use in this function." } */
+ ;
+ #pragma omp target uses_allocators (foo (traits_array), baz (traits_array))
+ ;
+ #pragma omp target uses_allocators (memspace(omp_no_such_space) : foo) /* { dg-error "'omp_no_such_space' undeclared" } */
+ ;
+ #pragma omp target uses_allocators (memspace(1) : foo) /* { dg-error "expected '\\\)' before numeric constant" } */
+ ;
+ #pragma omp target uses_allocators (memspace(omp_no_such_space) : foo, bar) /* { dg-error "memspace modifier must be constant enum of 'omp_memspace_handle_t' type" } */
+ ;
+ #pragma omp target uses_allocators (traits(xyz) : bar)
+ ;
+ #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space), traits(traits_array), memspace (omp_no_such_space) : bar) /* { dg-error "expected ':' before ',' token" } */
+ ;
+ #pragma omp target uses_allocators (traitz(traits_array), memspace(omp_high_bw_mem_space) : bar) uses_allocators(foo) /* { dg-error "'traitz' undeclared .first use in this function." } */
+ ;
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,53 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
+
+program main
+ use omp_lib
+ implicit none
+ integer, allocatable :: arr(:)
+ integer (omp_allocator_handle_kind) :: bar, foo
+
+ type (omp_alloctrait), parameter :: traits_array(*) = &
+ [omp_alloctrait(omp_atk_pinned,omp_atv_true),&
+ omp_alloctrait(omp_atk_partition,omp_atv_nearest)]
+
+ !$omp target allocate(bar : arr) uses_allocators(bar)
+ block
+ allocate(arr(100))
+ end block
+
+ !$omp target uses_allocators(omp_default_mem_alloc)
+ block
+ end block
+
+ !$omp target uses_allocators(bar(traits_array), foo (traits_array))
+ block
+ end block
+
+ !$omp target uses_allocators(traits(traits_array) : bar)
+ block
+ end block
+
+ !$omp target uses_allocators(memspace (omp_low_lat_mem_space) : bar)
+ block
+ end block
+
+ !$omp target uses_allocators(memspace (omp_high_bw_mem_space), traits(traits_array) : bar)
+ block
+ use iso_c_binding
+ type(c_ptr) :: ptr
+ integer(c_size_t) :: sz = 32
+ ptr = omp_alloc (sz, bar)
+ call omp_free (ptr, bar)
+ end block
+
+end program main
+
+! { dg-final { scan-tree-dump "pragma omp target allocate\\(allocator\\(bar\\):arr\\) private\\(bar\\) uses_allocators\\(bar: memspace\\(\\), traits\\(\\)\\)" "original" } }
+! { dg-final { scan-tree-dump "pragma omp target" "original" } }
+! { dg-final { scan-tree-dump "pragma omp target private\\(foo\\) uses_allocators\\(foo: memspace\\(\\), traits\\(traits_array\\)\\) private\\(bar\\) uses_allocators\\(bar: memspace\\(\\), traits\\(traits_array\\)\\)" "original" } }
+! { dg-final { scan-tree-dump "pragma omp target private\\(bar\\) uses_allocators\\(bar: memspace\\(\\), traits\\(traits_array\\)\\)" "original" } }
+! { dg-final { scan-tree-dump "pragma omp target private\\(bar\\) uses_allocators\\(bar: memspace\\(.\\), traits\\(\\)\\)" "original" } }
+! { dg-final { scan-tree-dump "pragma omp target private\\(bar\\) uses_allocators\\(bar: memspace\\(.\\), traits\\(traits_array\\)\\)" "original" } }
+! { dg-final { scan-tree-dump-times "__builtin_omp_init_allocator" 6 "gimple" } }
+! { dg-final { scan-tree-dump-times "__builtin_omp_destroy_allocator" 6 "gimple" } }
new file mode 100644
@@ -0,0 +1,44 @@
+! { dg-do compile }
+
+program main
+ use omp_lib
+ implicit none
+ integer (omp_allocator_handle_kind) :: bar, foo
+
+ type (omp_alloctrait), parameter :: traits_array(*) = &
+ [omp_alloctrait(omp_atk_pinned,omp_atv_true),&
+ omp_alloctrait(omp_atk_partition,omp_atv_nearest)]
+
+ !$omp target uses_allocators(omp_non_existant_alloc) ! { dg-error "'omp_non_existant_alloc' at .1. must be integer of 'omp_allocator_handle_kind' kind" }
+ block
+ end block
+
+ !$omp target uses_allocators(bar(traits_array), foo (traits_array), ) ! { dg-error "Expected name of allocator at .1." }
+ block
+ end block
+
+ !$omp target uses_allocators(traits(xyz) : bar) ! { dg-error "'xyz' at .1. must be of constant 'type.omp_alloctrait.' array type and have a constant initializer" }
+ block
+ end block
+
+ !$omp target uses_allocators(memspace(omp_non_existant_mem_space) : foo) ! { dg-error "'omp_non_existant_mem_space' at .1. is not a pre-defined memory space name" }
+ block
+ end block
+
+ !$omp target uses_allocators(traits(traits_array), traits(traits_array) : bar) ! { dg-error "Multiple traits modifiers at .1." }
+ block
+ end block
+
+ !$omp target uses_allocators(memspace(omp_default_mem_space), memspace(omp_default_mem_space) : foo) ! { dg-error "Multiple memspace modifiers at .1." }
+ block
+ end block
+
+ !$omp target uses_allocators(memspace(omp_default_mem_space), traits(traits_array), traits(traits_array) : foo) ! { dg-error "Only two modifiers are allowed on 'uses_allocators' clause at .1." }
+ block
+ end block
+
+ !$omp target uses_allocators(memspace(omp_default_mem_space), traits(traits_array) : foo, bar) ! { dg-error "When using modifiers, only a single allocator can be specified in each 'uses_allocators' clause at .1." }
+ block
+ end block
+
+end program main
new file mode 100644
@@ -0,0 +1,14 @@
+! { dg-do compile }
+
+program main
+ use omp_lib
+ implicit none
+ integer, allocatable :: arr(:)
+ integer (omp_allocator_handle_kind) :: bar
+
+ !$omp target allocate(bar : arr) ! { dg-error "allocator ''bar''in 'allocate' clause on target region is missing 'uses_allocators.bar.' clause" }
+ block
+ allocate(arr(100))
+ end block
+
+end program main
@@ -522,6 +522,9 @@ enum omp_clause_code {
/* OpenACC clause: nohost. */
OMP_CLAUSE_NOHOST,
+
+ /* OpenMP clause: uses_allocators. */
+ OMP_CLAUSE_USES_ALLOCATORS,
};
#undef DEFTREESTRUCT
@@ -769,6 +769,20 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
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(");
{
@@ -349,6 +349,7 @@ unsigned const char omp_clause_num_ops[] =
0, /* OMP_CLAUSE_IF_PRESENT */
0, /* OMP_CLAUSE_FINALIZE */
0, /* OMP_CLAUSE_NOHOST */
+ 3, /* OMP_CLAUSE_USES_ALLOCATORS */
};
const char * const omp_clause_code_name[] =
@@ -439,6 +440,7 @@ const char * const omp_clause_code_name[] =
"if_present",
"finalize",
"nohost",
+ "uses_allocators",
};
/* Unless specific to OpenACC, we tend to internally maintain OpenMP-centric
@@ -1824,6 +1824,15 @@ class auto_suppress_location_wrappers
#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)