@@ -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)
@@ -1873,6 +1873,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,
@@ -12922,6 +12922,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))
@@ -15651,6 +15653,213 @@ c_parser_omp_clause_allocate (c_parser *parser, tree list)
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 t = NULL_TREE, nl = list;
+ matching_parens parens;
+ if (!parens.require_open (parser))
+ return list;
+
+ tree memspace_expr = NULL_TREE;
+ tree traits_var = NULL_TREE;
+
+ struct item_tok
+ {
+ location_t loc;
+ tree id;
+ item_tok (void) : loc (UNKNOWN_LOCATION), id (NULL_TREE) {}
+ };
+ struct item { item_tok name, arg; };
+ auto_vec<item> *modifiers = NULL, *allocators = NULL;
+ auto_vec<item> *cur_list = new auto_vec<item> (4);
+
+ while (true)
+ {
+ item it;
+
+ if (c_parser_next_token_is (parser, CPP_NAME))
+ {
+ c_token *tok = c_parser_peek_token (parser);
+ it.name.id = tok->value;
+ it.name.loc = tok->location;
+ 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))
+ {
+ tok = c_parser_peek_token (parser);
+ it.arg.id = tok->value;
+ it.arg.loc = tok->location;
+ c_parser_consume_token (parser);
+ }
+ else
+ {
+ c_parser_error (parser, "expected identifier");
+ parens2.skip_until_found_close (parser);
+ goto end;
+ }
+ parens2.skip_until_found_close (parser);
+ }
+ }
+
+ cur_list->safe_push (it);
+
+ if (c_parser_next_token_is (parser, CPP_COMMA))
+ c_parser_consume_token (parser);
+ else if (c_parser_next_token_is (parser, CPP_COLON))
+ {
+ if (modifiers)
+ {
+ c_parser_error (parser, "expected %<)%>");
+ goto end;
+ }
+ else
+ {
+ c_parser_consume_token (parser);
+ modifiers = cur_list;
+ cur_list = new auto_vec<item> (4);
+ }
+ }
+ else if (c_parser_next_token_is (parser, CPP_CLOSE_PAREN))
+ {
+ gcc_assert (allocators == NULL);
+ allocators = cur_list;
+ cur_list = NULL;
+ break;
+ }
+ else
+ {
+ c_parser_error (parser, "expected %<)%>");
+ goto end;
+ }
+ }
+
+ if (modifiers)
+ for (unsigned i = 0; i < modifiers->length (); i++)
+ {
+ item& it = (*modifiers)[i];
+ const char *p = IDENTIFIER_POINTER (it.name.id);
+ int strcmp_traits = 1, strcmp_memspace = 1;
+
+ if ((strcmp_traits = strcmp ("traits", p)) == 0
+ || (strcmp_memspace = strcmp ("memspace", p)) == 0)
+ {
+ if ((strcmp_traits == 0 && traits_var != NULL_TREE)
+ || (strcmp_memspace == 0 && memspace_expr != NULL_TREE))
+ {
+ error_at (it.name.loc, "duplicate %qs modifier", p);
+ goto end;
+ }
+ t = lookup_name (it.arg.id);
+ if (t == NULL_TREE)
+ {
+ undeclared_variable (it.arg.loc, it.arg.id);
+ t = error_mark_node;
+ }
+ else if (strcmp_memspace == 0)
+ memspace_expr = t;
+ else if (strcmp_traits == 0)
+ traits_var = t;
+ else
+ gcc_unreachable ();
+ }
+ else
+ {
+ error_at (it.name.loc, "unknown modifier %qE", it.name.id);
+ goto end;
+ }
+ }
+
+ if (allocators)
+ {
+ if (modifiers)
+ {
+ if (allocators->length () > 1)
+ {
+ error_at ((*allocators)[1].name.loc,
+ "%<uses_allocators%> clause only accepts a single "
+ "allocator when using modifiers");
+ goto end;
+ }
+ else if ((*allocators)[0].arg.id)
+ {
+ error_at ((*allocators)[0].arg.loc,
+ "legacy %<%E(%E)%> traits syntax not allowed in "
+ "%<uses_allocators%> clause when using modifiers",
+ (*allocators)[0].name.id, (*allocators)[0].arg.id);
+ goto end;
+ }
+ }
+
+ for (unsigned i = 0; i < allocators->length (); i++)
+ {
+ item& it = (*allocators)[i];
+ t = lookup_name (it.name.id);
+ if (t == NULL_TREE)
+ {
+ undeclared_variable (it.name.loc, it.name.id);
+ goto end;
+ }
+ else if (t != error_mark_node)
+ {
+ tree t2 = NULL_TREE;
+ if (it.arg.id)
+ {
+ t2 = lookup_name (it.arg.id);
+ if (t2 == NULL_TREE)
+ {
+ undeclared_variable (it.arg.loc, it.arg.id);
+ goto end;
+ }
+ }
+ else
+ t2 = traits_var;
+
+ 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) = t2;
+ OMP_CLAUSE_CHAIN (c) = nl;
+ nl = c;
+ }
+ }
+ }
+ end:
+ if (cur_list)
+ delete cur_list;
+ if (modifiers)
+ delete modifiers;
+ if (allocators)
+ delete allocators;
+ parens.skip_until_found_close (parser);
+ return nl;
+}
+
/* OpenMP 4.0:
linear ( variable-list )
linear ( variable-list : expression )
@@ -17079,6 +17288,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";
@@ -21093,7 +21306,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,119 @@ 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 ((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 must be of %<omp_allocator_handle_t%> type");
+ remove = true;
+ break;
+ }
+ if (TREE_CODE (t) == CONST_DECL)
+ {
+ /* Currently for pre-defined allocators in libgomp, we do not
+ require additional init/fini inside target regions, so discard
+ such clauses. */
+ remove = true;
+
+ if (strcmp (IDENTIFIER_POINTER (DECL_NAME (t)),
+ "omp_null_allocator") == 0)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%<omp_null_allocator%> cannot be used in "
+ "%<uses_allocators%> clause");
+ break;
+ }
+
+ 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");
+ break;
+ }
+ }
+ 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;
+ break;
+ }
+ t = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c);
+ if (t != NULL_TREE)
+ {
+ bool type_err = false;
+
+ if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE
+ || DECL_SIZE (t) == NULL_TREE)
+ 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)
@@ -36649,6 +36649,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))
@@ -38901,6 +38903,234 @@ cp_parser_omp_clause_allocate (cp_parser *parser, tree list)
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 t = NULL_TREE, nl = list;
+ matching_parens parens;
+ if (!parens.require_open (parser))
+ return list;
+
+ tree memspace_expr = NULL_TREE;
+ tree traits_var = NULL_TREE;
+
+ struct item_tok
+ {
+ location_t loc;
+ tree id;
+ item_tok (void) : loc (UNKNOWN_LOCATION), id (NULL_TREE) {}
+ };
+ struct item { item_tok name, arg; };
+ auto_vec<item> *modifiers = NULL, *allocators = NULL;
+ auto_vec<item> *cur_list = new auto_vec<item> (4);
+
+ while (true)
+ {
+ item it;
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+ {
+ cp_token *tok = cp_lexer_peek_token (parser->lexer);
+ it.name.id = tok->u.value;
+ it.name.loc = 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);
+ it.arg.id = tok->u.value;
+ it.arg.loc = tok->location;
+ cp_lexer_consume_token (parser->lexer);
+ }
+ else
+ {
+ cp_parser_error (parser, "expected identifier");
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ goto end;
+ }
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/false,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ }
+ }
+
+ cur_list->safe_push (it);
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
+ cp_lexer_consume_token (parser->lexer);
+ else if (cp_lexer_next_token_is (parser->lexer, CPP_COLON))
+ {
+ if (modifiers)
+ {
+ cp_parser_error (parser, "expected %<)%>");
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ goto end;
+ }
+ else
+ {
+ cp_lexer_consume_token (parser->lexer);
+ modifiers = cur_list;
+ cur_list = new auto_vec<item> (4);
+ }
+ }
+ else if (cp_lexer_next_token_is (parser->lexer, CPP_CLOSE_PAREN))
+ {
+ gcc_assert (allocators == NULL);
+ allocators = cur_list;
+ cur_list = NULL;
+ break;
+ }
+ else
+ {
+ cp_parser_error (parser, "expected %<)%>");
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ goto end;
+ }
+ }
+
+ if (modifiers)
+ for (unsigned i = 0; i < modifiers->length (); i++)
+ {
+ item& it = (*modifiers)[i];
+ const char *p = IDENTIFIER_POINTER (it.name.id);
+ int strcmp_traits = 1, strcmp_memspace = 1;
+
+ if ((strcmp_traits = strcmp ("traits", p)) == 0
+ || (strcmp_memspace = strcmp ("memspace", p)) == 0)
+ {
+ if ((strcmp_traits == 0 && traits_var != NULL_TREE)
+ || (strcmp_memspace == 0 && memspace_expr != NULL_TREE))
+ {
+ error_at (it.name.loc, "duplicate %qs modifier", p);
+ goto end;
+ }
+ t = cp_parser_lookup_name_simple (parser, it.arg.id, it.arg.loc);
+ if (t == error_mark_node)
+ {
+ cp_parser_name_lookup_error (parser, it.arg.id, t, NLE_NULL,
+ it.arg.loc);
+ }
+ else if (strcmp_memspace == 0)
+ memspace_expr = t;
+ else if (strcmp_traits == 0)
+ traits_var = t;
+ else
+ gcc_unreachable ();
+ }
+ else
+ {
+ error_at (it.name.loc, "unknown modifier %qE", it.name.id);
+ goto end;
+ }
+ }
+
+ if (allocators)
+ {
+ if (modifiers)
+ {
+ if (allocators->length () > 1)
+ {
+ error_at ((*allocators)[1].name.loc,
+ "%<uses_allocators%> clause only accepts a single "
+ "allocator when using modifiers");
+ goto end;
+ }
+ else if ((*allocators)[0].arg.id)
+ {
+ error_at ((*allocators)[0].arg.loc,
+ "legacy %<%E(%E)%> traits syntax not allowed in "
+ "%<uses_allocators%> clause when using modifiers",
+ (*allocators)[0].name.id, (*allocators)[0].arg.id);
+ goto end;
+ }
+ }
+
+ for (unsigned i = 0; i < allocators->length (); i++)
+ {
+ item& it = (*allocators)[i];
+ t = cp_parser_lookup_name_simple (parser, it.name.id, it.name.loc);
+ if (t == error_mark_node)
+ {
+ cp_parser_name_lookup_error (parser, it.name.id, t, NLE_NULL,
+ it.name.loc);
+ goto end;
+ }
+ else if (t != error_mark_node)
+ {
+ tree t2 = NULL_TREE;
+ if (it.arg.id)
+ {
+ t2 = cp_parser_lookup_name_simple (parser, it.arg.id,
+ it.arg.loc);
+ if (t2 == error_mark_node)
+ {
+ cp_parser_name_lookup_error (parser, it.arg.id, t2,
+ NLE_NULL, it.arg.loc);
+ goto end;
+ }
+ }
+ else
+ t2 = traits_var;
+
+ 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) = t2;
+ OMP_CLAUSE_CHAIN (c) = nl;
+ nl = c;
+ }
+ }
+ }
+ end:
+ if (cur_list)
+ delete cur_list;
+ if (modifiers)
+ delete modifiers;
+ if (allocators)
+ delete allocators;
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/false,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return nl;
+}
+
/* OpenMP 2.5:
lastprivate ( variable-list )
@@ -40453,6 +40683,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;
@@ -44464,7 +44698,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,
@@ -7766,6 +7766,113 @@ 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 (t) == FIELD_DECL)
+ {
+ sorry_at (OMP_CLAUSE_LOCATION (c), "class members not yet "
+ "supported in %<uses_allocators%> clause");
+ remove = true;
+ break;
+ }
+ t = convert_from_reference (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;
+ break;
+ }
+ if (TREE_CODE (t) == CONST_DECL)
+ {
+ /* Currently for pre-defined allocators in libgomp, we do not
+ require additional init/fini inside target regions, so discard
+ such clauses. */
+ remove = true;
+
+ if (strcmp (IDENTIFIER_POINTER (DECL_NAME (t)),
+ "omp_null_allocator") == 0)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%<omp_null_allocator%> cannot be used in "
+ "%<uses_allocators%> clause");
+ break;
+ }
+
+ 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");
+ break;
+ }
+ }
+ 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;
+ break;
+ }
+ t = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c);
+ if (t != NULL_TREE)
+ {
+ bool type_err = false;
+
+ if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE
+ || DECL_SIZE (t) == NULL_TREE)
+ 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
+ {
+ 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)
@@ -1424,6 +1424,20 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n)
case OMP_LINEAR_UVAL: fputs ("uval(", dumpfile); break;
default: break;
}
+ else if (list_type == OMP_LIST_USES_ALLOCATORS)
+ {
+ show_symbol (n->sym);
+ fputs ("(memspace:", dumpfile);
+ if (n->memspace_sym)
+ show_symbol (n->traits_sym);
+ fputs (",traits:", dumpfile);
+ if (n->memspace_sym)
+ show_symbol (n->traits_sym);
+ fputc (')', dumpfile);
+ if (n->next)
+ fputc (',', dumpfile);
+ continue;
+ }
fprintf (dumpfile, "%s", n->sym ? n->sym->name : "omp_all_memory");
if (list_type == OMP_LIST_LINEAR && n->u.linear_op != OMP_LINEAR_DEFAULT)
fputc (')', dumpfile);
@@ -1690,6 +1704,7 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses)
case OMP_LIST_ALLOCATE: type = "ALLOCATE"; break;
case OMP_LIST_SCAN_IN: type = "INCLUSIVE"; break;
case OMP_LIST_SCAN_EX: type = "EXCLUSIVE"; break;
+ case OMP_LIST_USES_ALLOCATORS: type = "USES_ALLOCATORS"; break;
default:
gcc_unreachable ();
}
@@ -1354,6 +1354,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;
}
@@ -1395,6 +1397,7 @@ enum
OMP_LIST_NONTEMPORAL,
OMP_LIST_ALLOCATE,
OMP_LIST_HAS_DEVICE_ADDR,
+ OMP_LIST_USES_ALLOCATORS,
OMP_LIST_NUM /* Must be the last. */
};
@@ -986,6 +986,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
};
@@ -1402,6 +1403,528 @@ gfc_match_omp_clause_reduction (char pc, gfc_omp_clauses *c, bool openacc,
return MATCH_YES;
}
+/* 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 match
+gfc_match_omp_clause_uses_allocators (gfc_omp_clauses *c)
+{
+ char buffer[GFC_MAX_SYMBOL_LEN + 1];
+ gfc_symbol *sym;
+ gfc_symbol *memspace_sym= NULL;
+ gfc_symbol *traits_sym= NULL;
+ locus traits_sym_loc;
+ match m, ret = MATCH_ERROR;
+
+ if (gfc_match ("uses_allocators ( ") != MATCH_YES)
+ return MATCH_NO;
+
+ struct item_tok
+ {
+ locus loc;
+ char *str;
+ item_tok (void) : str (NULL) {}
+ ~item_tok (void) { if (str) free (str); }
+ };
+ struct item { item_tok name, arg; };
+ auto_vec<item> *modifiers = NULL, *allocators = NULL;
+ auto_vec<item> *cur_list = new auto_vec<item> (4);
+
+ gfc_symbol *allocator_handle_kind;
+
+ if (gfc_find_symbol ("omp_allocator_handle_kind", NULL, 1, &sym)
+ || sym == NULL
+ || sym->attr.dimension
+ || sym->value == NULL
+ || 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;
+
+ while (true)
+ {
+ item it;
+
+ m = gfc_match_name (buffer);
+ if (m == MATCH_YES)
+ {
+ it.name.str = xstrdup (buffer);
+ it.name.loc = gfc_current_locus;
+ }
+ else
+ {
+ gfc_error ("Expected identifier at %C");
+ goto error;
+ }
+
+ if (gfc_match_char ('(') == MATCH_YES)
+ {
+ m = gfc_match_name (buffer);
+ if (m == MATCH_YES)
+ {
+ it.arg.str = xstrdup (buffer);
+ it.arg.loc = gfc_current_locus;
+ }
+ else
+ {
+ gfc_error ("Expected identifier at %C");
+ goto error;
+ }
+ if (gfc_match_char (')') != MATCH_YES)
+ {
+ gfc_error ("Expected %<)%> at %C");
+ goto error;
+ }
+ }
+
+ cur_list->safe_push (it);
+ it.name.str = NULL;
+ it.arg.str = NULL;
+
+ if (gfc_match (" , ") == MATCH_YES)
+ continue;
+ else if (gfc_match (" : ") == MATCH_YES)
+ {
+ if (modifiers)
+ {
+ gfc_error ("expected %<)%> at %C");
+ goto error;
+ }
+ else
+ {
+ modifiers = cur_list;
+ cur_list = new auto_vec<item> (4);
+ }
+ }
+ else if (gfc_match_char (')') == MATCH_YES)
+ {
+ gcc_assert (allocators == NULL);
+ allocators = cur_list;
+ cur_list = NULL;
+ break;
+ }
+ else
+ {
+ gfc_error ("expected %<)%> at %C");
+ goto error;
+ }
+ }
+
+ if (modifiers)
+ for (unsigned i = 0; i < modifiers->length (); i++)
+ {
+ item& it = (*modifiers)[i];
+ const char *p = it.name.str;
+ int strcmp_traits = 1, strcmp_memspace = 1;
+ gfc_symbol *sym;
+
+ if ((strcmp_traits = strcmp ("traits", p)) == 0
+ || (strcmp_memspace = strcmp ("memspace", p)) == 0)
+ {
+ if ((strcmp_traits == 0 && traits_sym != NULL)
+ || (strcmp_memspace == 0 && memspace_sym != NULL))
+ {
+ gfc_error ("duplicate %qs modifier at %L", p, &it.name.loc);
+ goto error;
+ }
+ if (gfc_find_symbol (it.arg.str, NULL, 1, &sym) || sym == NULL)
+ {
+ gfc_error ("Symbol %qs at %L is ambiguous",
+ it.arg.str, &it.arg.loc);
+ goto error;
+ }
+ else if (strcmp_memspace == 0)
+ {
+ memspace_sym = sym;
+
+ /* We have a memspace specified, now check if it is valid.
+ Start with finding if we have the standards specified
+ 'omp_memspace_handle_kind' available. */
+ if (gfc_find_symbol ("omp_memspace_handle_kind", NULL, 1, &sym)
+ || sym == NULL
+ || sym->attr.dimension
+ || sym->value == NULL
+ || 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 %L",
+ &it.arg.loc);
+ goto error;
+ }
+
+ gfc_symbol *memspace_handle_kind = sym;
+
+ 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. */
+ || !startswith (memspace_sym->name, "omp_")
+ || !endswith (memspace_sym->name, "_mem_space"))
+ {
+ gfc_error ("%<%s%> at %L is not a pre-defined memory space "
+ "name", memspace_sym->name, &it.arg.loc);
+ goto error;
+ }
+ }
+ else if (strcmp_traits == 0)
+ {
+ traits_sym = sym;
+ traits_sym_loc = it.arg.loc;
+ }
+ else
+ gcc_unreachable ();
+ }
+ else
+ {
+ gfc_error ("unknown modifier %qs at %L", p, &it.name.loc);
+ goto error;
+ }
+ }
+
+ if (allocators)
+ {
+ if (modifiers)
+ {
+ if (allocators->length () > 1)
+ {
+ gfc_error ("%<uses_allocators%> clause only accepts a single "
+ "allocator when using modifiers at %L",
+ &(*allocators)[1].name.loc);
+ goto error;
+ }
+ else if ((*allocators)[0].arg.str)
+ {
+ gfc_error ("legacy %<%s(%s)%> traits syntax not allowed in "
+ "%<uses_allocators%> clause when using modifiers at %L",
+ (*allocators)[0].name.str, (*allocators)[0].arg.str,
+ &(*allocators)[0].arg.loc);
+ goto error;
+ }
+ }
+
+ for (unsigned i = 0; i < allocators->length (); i++)
+ {
+ item& it = (*allocators)[i];
+
+ gfc_symbol *allocator_sym;
+ locus allocator_sym_loc;
+
+ if (gfc_find_symbol (it.name.str, NULL, 1, &allocator_sym) != 0
+ || allocator_sym == NULL)
+ {
+ gfc_error ("Symbol %qs at %L is ambiguous",
+ it.name.str, &it.name.loc);
+ goto error;
+ }
+ allocator_sym_loc = it.name.loc;
+
+ gfc_symbol *curr_traits_sym;
+ locus curr_traits_sym_loc;
+
+ if (it.arg.str)
+ {
+ if (gfc_find_symbol (it.arg.str, NULL, 1, &curr_traits_sym)
+ || curr_traits_sym == NULL)
+ {
+ gfc_error ("Symbol %qs at %L is ambiguous",
+ it.arg.str, &it.arg.loc);
+ goto error;
+ }
+ curr_traits_sym_loc = it.arg.loc;
+ }
+ else
+ {
+ curr_traits_sym = traits_sym;
+ curr_traits_sym_loc = traits_sym_loc;
+ }
+
+ if (curr_traits_sym)
+ {
+ if (curr_traits_sym->ts.type != BT_DERIVED
+ || strcmp (curr_traits_sym->ts.u.derived->name,
+ "omp_alloctrait") != 0
+ || curr_traits_sym->attr.flavor != FL_PARAMETER
+ || curr_traits_sym->as->rank != 1)
+ {
+ gfc_error ("%<%s%> at %L must be of constant "
+ "%<type(omp_alloctrait)%> array type and have a "
+ "constant initializer", curr_traits_sym->name,
+ &curr_traits_sym_loc);
+ goto error;
+ }
+ gfc_set_sym_referenced (curr_traits_sym);
+ }
+
+ 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 %L must be integer of %<%s%> kind",
+ allocator_sym->name, &allocator_sym_loc,
+ allocator_handle_kind->name);
+ goto error;
+ }
+
+ if (allocator_sym->attr.flavor == FL_PARAMETER)
+ {
+ if (strcmp (allocator_sym->name, "omp_null_allocator") == 0)
+ {
+ gfc_error ("%<omp_null_allocator%> cannot be used in "
+ "%<uses_allocators%> clause at %L",
+ &allocator_sym_loc);
+ goto error;
+ }
+
+ /* Check if identifier is a 'omp_..._mem_alloc' pre-defined
+ allocator. */
+ if (!startswith (allocator_sym->name, "omp_")
+ || !endswith (allocator_sym->name, "_mem_alloc"))
+ {
+ gfc_error ("%<%s%> at %L is not a pre-defined memory "
+ "allocator", allocator_sym->name,
+ &allocator_sym_loc);
+ 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 = curr_traits_sym;
+ n->where = it.name.loc;
+
+ n->next = c->lists[OMP_LIST_USES_ALLOCATORS];
+ c->lists[OMP_LIST_USES_ALLOCATORS] = n;
+ }
+ }
+ }
+
+ ret = MATCH_YES;
+
+ end:
+ if (cur_list)
+ delete cur_list;
+ if (modifiers)
+ delete modifiers;
+ if (allocators)
+ delete allocators;
+ return ret;
+
+ error:
+ ret = MATCH_ERROR;
+ gfc_error_check ();
+ goto end;
+
+#if 0
+ 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 (deprecated) traits "
+ "array list 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)
+ {
+ 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)
+ {
+ if (gfc_find_symbol ("omp_memspace_handle_kind", NULL, 1, &sym)
+ || sym == NULL
+ || sym->attr.dimension
+ || sym->value == NULL
+ || 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;
+ }
+ gfc_symbol *memspace_handle_kind = sym;
+
+ 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. */
+ || !startswith (memspace_sym->name, "omp_")
+ || !endswith (memspace_sym->name, "_mem_space"))
+ {
+ 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)
+ {
+ /* Check if identifier is a 'omp_..._mem_alloc' pre-defined
+ allocator. */
+ if (!startswith (allocator_sym->name, "omp_")
+ || !endswith (allocator_sym->name, "_mem_alloc"))
+ {
+ 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;
+#endif
+}
/* Match with duplicate check. Matches 'name'. If expr != NULL, it
then matches '(expr)', otherwise, if open_parens is true,
@@ -2971,6 +3494,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
@@ -3697,7 +4223,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)
@@ -6330,7 +6856,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,6 @@ 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;
if (mpz_cmp_si (c->repeat, 1) > 0)
{
@@ -6371,6 +6367,11 @@ gfc_conv_array_initializer (tree type, gfc_expr * expr)
else
range = NULL;
+ if (range == NULL || mpz_cmp_si (c->offset, 0) != 0)
+ index = gfc_conv_mpz_to_tree (c->offset, gfc_index_integer_kind);
+ else
+ index = NULL_TREE;
+
gfc_init_se (&se, NULL);
switch (c->expr->expr_type)
{
@@ -2719,9 +2719,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);
@@ -3701,6 +3708,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;
}
@@ -6126,6 +6156,8 @@ gfc_split_omp_clauses (gfc_code *code,
= code->ext.omp_clauses->device;
clausesa[GFC_OMP_SPLIT_TARGET].thread_limit
= code->ext.omp_clauses->thread_limit;
+ clausesa[GFC_OMP_SPLIT_TARGET].lists[OMP_LIST_USES_ALLOCATORS]
+ = code->ext.omp_clauses->lists[OMP_LIST_USES_ALLOCATORS];
for (int i = 0; i < OMP_DEFAULTMAP_CAT_NUM; i++)
clausesa[GFC_OMP_SPLIT_TARGET].defaultmap[i]
= code->ext.omp_clauses->defaultmap[i];
@@ -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)
@@ -10831,6 +10831,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_CLAUSE_BIND:
case OMP_CLAUSE_IF_PRESENT:
case OMP_CLAUSE_FINALIZE:
+ case OMP_CLAUSE_USES_ALLOCATORS:
break;
case OMP_CLAUSE_ORDER:
@@ -10945,6 +10946,49 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
remove = true;
break;
}
+ if ((omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0
+ && OMP_CLAUSE_ALLOCATE_ALLOCATOR (c)
+ && TREE_CODE (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c)) != INTEGER_CST)
+ {
+ tree allocator = OMP_CLAUSE_ALLOCATE_ALLOCATOR (c);
+ tree clauses = NULL_TREE;
+
+ /* Get clause list of the nearest enclosing target construct. */
+ if (ctx->code == OMP_TARGET)
+ clauses = *orig_list_p;
+ else
+ {
+ struct gimplify_omp_ctx *tctx = ctx->outer_context;
+ while (tctx && tctx->code != OMP_TARGET)
+ tctx = tctx->outer_context;
+ if (tctx)
+ clauses = tctx->clauses;
+ }
+
+ if (clauses)
+ {
+ tree uc;
+ if (TREE_CODE (allocator) == MEM_REF
+ || TREE_CODE (allocator) == INDIRECT_REF)
+ allocator = TREE_OPERAND (allocator, 0);
+ for (uc = clauses; uc; uc = OMP_CLAUSE_CHAIN (uc))
+ if (OMP_CLAUSE_CODE (uc) == OMP_CLAUSE_USES_ALLOCATORS)
+ {
+ tree uc_allocator
+ = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (uc);
+ if (operand_equal_p (allocator, uc_allocator))
+ break;
+ }
+ if (uc == NULL_TREE)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "allocator %<%qE%> "
+ "requires %<uses_allocators(%E)%> clause in "
+ "target region", allocator, allocator);
+ remove = true;
+ break;
+ }
+ }
+ }
if (gimplify_expr (&OMP_CLAUSE_ALLOCATE_ALLOCATOR (c), pre_p, NULL,
is_gimple_val, fb_rvalue) == GS_ERROR)
{
@@ -14271,6 +14315,73 @@ 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);
+
+ 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)
@@ -4849,7 +4849,12 @@ lower_private_allocate (tree var, tree new_var, tree &allocator,
allocator = TREE_PURPOSE (allocator);
}
if (TREE_CODE (allocator) != INTEGER_CST)
- allocator = build_outer_var_ref (allocator, ctx);
+ {
+ if (is_task_ctx (ctx))
+ allocator = build_receiver_ref (allocator, false, ctx);
+ else
+ allocator = build_outer_var_ref (allocator, ctx);
+ }
allocator = fold_convert (pointer_sized_int_node, allocator);
if (TREE_CODE (allocator) != INTEGER_CST)
{
@@ -5833,7 +5838,7 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist,
if (TREE_CODE (allocator) == TREE_LIST)
allocator = TREE_PURPOSE (allocator);
if (TREE_CODE (allocator) != INTEGER_CST)
- allocator = build_outer_var_ref (allocator, ctx);
+ allocator = build_receiver_ref (allocator, false, ctx);
allocator = fold_convert (pointer_sized_int_node,
allocator);
allocate_ptr = unshare_expr (x);
@@ -6153,7 +6158,7 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist,
if (TREE_CODE (allocator) == TREE_LIST)
allocator = TREE_PURPOSE (allocator);
if (TREE_CODE (allocator) != INTEGER_CST)
- allocator = build_outer_var_ref (allocator, ctx);
+ allocator = build_receiver_ref (allocator, false, ctx);
allocator = fold_convert (pointer_sized_int_node,
allocator);
allocate_ptr = unshare_expr (x);
@@ -12223,6 +12228,8 @@ create_task_copyfn (gomp_task *task_stmt, omp_context *ctx)
allocator = *tcctx.cb.decl_map->get (allocator);
tree a = build_simple_mem_ref_loc (loc, sarg);
allocator = omp_build_component_ref (a, allocator);
+ if (POINTER_TYPE_P (TREE_TYPE (allocator)))
+ allocator = build_simple_mem_ref (allocator);
}
allocator = fold_convert (pointer_sized_int_node, allocator);
tree a = builtin_decl_explicit (BUILT_IN_GOMP_ALLOC);
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 parallel uses_allocators (memspace(omp_high_bw_mem_space), traits(foo_traits) : bar)
+ ;
+ #pragma omp target parallel 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\\(bar\\) uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\), traits\\(foo_traits\\)\\) private\\(foo\\) uses_allocators\\(foo: memspace\\(\\), 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,39 @@
+/* { 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." "" { 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 (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' has not been declared" "" { target c++ } .-1 } */
+ #pragma omp target uses_allocators (memspace(1) : foo) /* { dg-error "expected identifier before numeric constant" } */
+ ; /* { dg-error "expected '\\\)' before ':' token" "" { target c } .-1 } */
+ #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 "'omp_no_such_space' has not been declared" "" { target c++ } .-1 } */
+ #pragma omp target uses_allocators (traits(xyz) : bar) /* { dg-error "traits array must be of 'const omp_alloctrait_t \\\[\\\]' type" "" { target c } } */
+ ; /* { dg-error "'xyz' has not been declared" "" { target c++ } .-1 } */
+ #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" } */
+ ;
+ #pragma omp target uses_allocators (traitz(traits_array), memspace(omp_high_bw_mem_space) : bar) /* { dg-error "unknown modifier 'traitz'" } */
+ ;
+ #pragma omp target uses_allocators (omp_null_allocator) /* { dg-error "'omp_null_allocator' cannot be used in 'uses_allocators' clause" } */
+ ;
+ #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 "legacy 'foo\\\(foo_traits\\\)' traits syntax not allowed in 'uses_allocators' clause when using modifiers" } */
+ ;
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,31 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */
+
+#include <omp.h>
+
+int main (void)
+{
+ omp_allocator_handle_t memspace, traits;
+ const omp_alloctrait_t mytraits[] = { { omp_atk_pinned, omp_atv_true },
+ { omp_atk_partition, omp_atv_nearest } };
+ #pragma omp target uses_allocators (memspace)
+ ;
+ #pragma omp target uses_allocators (traits)
+ ;
+ #pragma omp target uses_allocators (traits, memspace)
+ ;
+ #pragma omp target uses_allocators (traits (mytraits))
+ ;
+ #pragma omp target uses_allocators (memspace (mytraits), omp_default_mem_alloc)
+ ;
+ return 0;
+}
+
+/* { dg-final { scan-tree-dump "pragma omp target private\\(memspace\\) uses_allocators\\(memspace: memspace\\(\\), traits\\(\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target private\\(traits\\) uses_allocators\\(traits: memspace\\(\\), traits\\(\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target private\\(memspace\\) uses_allocators\\(memspace: memspace\\(\\), traits\\(\\)\\) private\\(traits\\) uses_allocators\\(traits: memspace\\(\\), traits\\(\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target private\\(traits\\) uses_allocators\\(traits: memspace\\(\\), traits\\(mytraits\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target private\\(memspace\\) uses_allocators\\(memspace: memspace\\(\\), traits\\(mytraits\\)\\)" "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" } } */
@@ -80,7 +80,8 @@ subroutine foo(x, y)
!$omp target teams distribute parallel do private (x) firstprivate (y) &
!$omp allocate ((omp_default_mem_alloc + 0):z) allocate &
- !$omp (omp_default_mem_alloc: x, y) allocate (h: r) lastprivate (z) reduction(+:r)
+ !$omp (omp_default_mem_alloc: x, y) allocate (h: r) lastprivate (z) reduction(+:r) &
+ !$omp uses_allocators (h)
do i = 1, 10
call bar (0, x, z);
call bar2 (1, y, r);
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 parallel uses_allocators(memspace (omp_low_lat_mem_space) : bar)
+ block
+ end block
+
+ !$omp target parallel 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,52 @@
+! { 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 "Symbol 'omp_non_existant_alloc' at .1. is ambiguous" }
+ block
+ end block
+
+ !$omp target uses_allocators(bar(traits_array), foo (traits_array), ) ! { dg-error "Expected identifier at .1." }
+ block ! { dg-error "Failed to match clause at .1." "" { target *-*-* } .-1 }
+ end block
+
+ !$omp target uses_allocators(traits(xyz) : bar) ! { dg-error "Symbol 'xyz' at .1. is ambiguous" }
+ block
+ end block
+
+ !$omp target uses_allocators(memspace(omp_non_existant_mem_space) : foo) ! { dg-error "Symbol 'omp_non_existant_mem_space' at .1. is ambiguous" }
+ block
+ end block
+
+ !$omp target uses_allocators(traits(traits_array), traits(traits_array) : bar) ! { dg-error "duplicate 'traits' modifier at .1." }
+ block
+ end block
+
+ !$omp target uses_allocators(memspace(omp_default_mem_space), memspace(omp_default_mem_space) : foo) ! { dg-error "duplicate 'memspace' modifier at .1." }
+ block
+ end block
+
+ !$omp target uses_allocators(memspace(omp_default_mem_space), traits(traits_array), traits(traits_array) : foo) ! { dg-error "duplicate 'traits' modifier at .1." }
+ block
+ end block
+
+ !$omp target uses_allocators (omp_null_allocator) ! { dg-error "'omp_null_allocator' cannot be used in 'uses_allocators' clause at .1." }
+ block
+ end block
+
+ !$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 at .1." }
+ block
+ end block
+
+ !$omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo(foo_traits)) ! { dg-error "legacy 'foo\\\(foo_traits\\\)' traits syntax not allowed in 'uses_allocators' clause when using modifiers 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'' requires 'uses_allocators.bar.' clause in target region" }
+ 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
@@ -1874,6 +1874,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)