On 2022/5/7 12:40 AM, Tobias Burnus wrote:
Can please also handle the new clause in Fortran's dump-parse-tree.cc?
I did see some split handling in C, but not in Fortran; do you also need
to up update gfc_split_omp_clauses in Fortran's trans-openmp.cc?
Done.
Actually, glancing at the testcases, no combined construct (like
"omp target parallel") is used, I think that would be useful because of ↑.
Okay, added some to testcases.
+/* OpenMP 5.2:
+ uses_allocators ( allocator-list )
That's not completely true: uses_allocators is OpenMP 5.1.
However, 5.1 only supports (for non-predefined allocators):
uses_allocators( allocator(traits) )
while OpenMP 5.2 added modifiers:
uses_allocatrors( traits(...), memspace(...) : allocator )
and deprecated the 5.1 'allocator(traits)'. (Scheduled for removal in OMP 6.0)
The advantage of 5.2 syntax is that a memory space can be defined.
I supported both syntaxes, that's why I designated it as "5.2".
BTW: This makes uses_allocators the first OpenMP 5.2 feature which
will make it into GCC :-)
:)
gcc/fortran/openmp.cc:
+ 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;
I think you rather want to use
gfc_find_symbol ("omp_...", NULL, true, &sym)
|| sym == NULL
where true is for parent_flag to search also the parent namespace.
(The function returns 1 if the symbol is ambiguous, 0 otherwise -
including 0 + sym == NULL when the symbol could not be found.)
|| sym->attr.flavor != FL_PARAMETER
|| sym->ts.type != BT_INTEGER
|| sym->attr.dimension
Looks cleaner than to access sym->value. The attr.dimension is just
to makes sure the user did not smuggle an array into this.
(Invalid as omp_... is a reserved namespace but users will still do
this and some are good in finding ICE as hobby.)
Well, the intention here is to search for "omp_allocator_handle_kind" and
"omp_memspace_handle_kind",
and use their value to check if the kinds are the same as declared allocator
handles and memspace constant.
Not to generally search for "omp_...".
However the sym->attr.dimension test seems useful, added in new v2 patch.
However, I fear that will fail for the following two examples (both untested):
use omp_lib, my_kind = omp_allocator_handle_kind
integer(my_kind) :: my_allocator
as this gives 'my_kind' in the symtree->name (while symtree->n.sym->name is
"omp_...").
Hence, by searching the symtree for 'omp_...' the symbol will not be found.
It will likely also fail for the following more realistic example:
...
subroutine foo
use m
use omp_lib, only: omp_alloctrait
...
!$omp target uses_allocators(my_allocator(traits_array)
allocate(my_allocator:A) firstprivate(A)
...
!$omp end target
end
If someone wants to use OpenMP allocators, but intentionally only imports
insufficient standard symbols from omp_lib,
then he/she is on their own :)
The specification really makes this quite clear: omp_allocator_handle_kind,
omp_alloctrait, omp_memspace_handle_kind are
all part of the same package.
In this case, omp_allocator_handle_kind is not in the namespace of 'foo'
but the code should be still valid. Thus, an alternative would be to hard-code
the value - as done for the depobj. As we have:
integer, parameter :: omp_allocator_handle_kind = c_intptr_t
integer, parameter :: omp_memspace_handle_kind = c_intptr_t
that would be
sym->ts.type == BT_CHARACTER
sym->ts.kind == gfc_index_integer_kind
for the allocator variable and the the memspace kind.
However, I grant that either example is not very typical. The second one is more
natural – such a code will very likely be written in the real world. But not
with uses_allocators but rather with "!$omp requires dynamic_allocators" and
omp_init_allocator().
Thoughts?
As above. I mean, what is so hard with including "use omp_lib" where you need
it? :D
* * *
gcc/fortran/openmp.cc
+ if (++i > 2)
+ {
+ gfc_error ("Only two modifiers are allowed on %<uses_allocators%> "
+ "clause at %C");
+ goto error;
+ }
+
Is this really needed? There is a check for multiple traits and multiple
memspace
Thus, 'trait(),memspace(),trait()' is already handled and
'trait(),something' give a break and will lead to an error as in that case
a ':' and not ',something' is expected.
I think it could be worth reminding that limitation, instead of a generic error.
+ if (gfc_match_char ('(') == MATCH_YES)
+ {
+ if (memspace_seen || traits_seen)
+ {
+ gfc_error ("Modifiers cannot be used with legacy "
+ "array syntax at %C");
I wouldn't uses the term 'array synax' to denote
uses_allocators(allocator (alloc_array) )
How about:
error: "Using both modifiers and allocator variable with traits argument"
(And I think 'deprecated' is better than 'legacy', if we really want to use it.)
I've changed it to "(deprecated) traits array list syntax", is that better?
+ 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))
I think the gfc_is_constant_expr is unreachable as you already
have checked FL_PARAMETER. Thus, you can remove the last two
lines.
Okay.
[Regarding the traits_sym->ts.u.derived->name, I am not sure whether that
won't fail with
use omp_lib, trait_t => omp_alloctrait
but I have not checked. It likely does work correctly.]
+ /* 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')
I wonder whether that's not more readable written as:
|| !startswith (memspace_sym->name, "omp_")
|| !endswith (memspace_sym->name, "_mem_space")
Thanks, didn't know it was this convenient :)
I've attached v2 of the patch. Currently in testing.
Thanks,
Chung-Lin
diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index 3a7cecdf087..be3e6ff697e 100644
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
@@ -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)
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 777cdc65572..5066e137cf4 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -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:
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index 54864c2ec41..7f8944f81d6 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -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,
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 129dd727ef3..bbdec92780b 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -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,233 @@ 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 t = NULL_TREE, nl = list;
+ 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);
+ 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);
+ }
+
+ if (!parens2.require_close (parser))
+ {
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+
+ 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 (t == error_mark_node)
+ return list;
+ }
+ has_modifiers = true;
+ }
+ }
+
+ 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 if (c_parser_next_token_is (parser, CPP_COMMA))
+ c_parser_error (parser, "modifiers can only be used with "
+ "a single allocator in %<uses_allocators%> "
+ "clause");
+ }
+ else
+ c_parser_error (parser, "expected identifier");
+ }
+ else
+ {
+ 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);
+ t = error_mark_node;
+ }
+ c_parser_consume_token (parser);
+
+ traits_var = NULL_TREE;
+ 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);
+ traits_var = error_mark_node;
+ }
+ c_parser_consume_token (parser);
+ }
+ else
+ c_parser_error (parser, "expected identifier");
+ parens2.require_close (parser);
+ }
+
+ if (t != error_mark_node && traits_var != 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) = 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 +17279,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 +21294,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)
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index e130196a3a7..0e1f33b655d 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -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)
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 2235da10c7c..e041bc669a9 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -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,247 @@ 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 = list;
+ 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)
+ {
+ 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 (!parens2.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;
+ }
+ }
+ has_modifiers = true;
+ }
+ }
+
+ 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 if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
+ cp_parser_error (parser, "modifiers can only be used with "
+ "a single allocator in %<uses_allocators%> "
+ "clause");
+ }
+ else
+ cp_parser_error (parser, "expected identifier");
+ }
+ else
+ {
+ 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);
+
+ traits_var = NULL_TREE;
+ 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);
+ cp_lexer_consume_token (parser->lexer);
+ }
+ else
+ cp_parser_error (parser, "expected identifier");
+ parens2.require_close (parser);
+ }
+
+ if (t != error_mark_node && traits_var != 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) = 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);
+ }
+ }
+
+ cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/false,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return nl;
+}
+
/* OpenMP 2.5:
lastprivate ( variable-list )
@@ -40283,6 +40526,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 +44538,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,
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 377f61113c0..c4ff73e7899 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -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 %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)
diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc
index 3635460bffd..3ac7fc846ac 100644
--- a/gcc/fortran/dump-parse-tree.cc
+++ b/gcc/fortran/dump-parse-tree.cc
@@ -1423,6 +1423,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->name);
if (list_type == OMP_LIST_LINEAR && n->u.linear_op != OMP_LINEAR_DEFAULT)
fputc (')', dumpfile);
@@ -1689,6 +1703,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 ();
}
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 7bf1d5a0452..18e685ca1b1 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -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. */
};
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 714148138c2..a187e75e1fe 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -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,234 @@ 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_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;
+
+ 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;
+ }
+ 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 (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 (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;
+
+ 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 +3153,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 +3882,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 +6514,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)
diff --git a/gcc/fortran/trans-array.cc b/gcc/fortran/trans-array.cc
index 05134952db4..a2a2b889d03 100644
--- a/gcc/fortran/trans-array.cc
+++ b/gcc/fortran/trans-array.cc
@@ -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)
{
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 43d59abe9e0..b094b17f054 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -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;
}
@@ -6074,6 +6104,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];
diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def
index cd79ad45167..18a1bec8724 100644
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
@@ -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)
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 2588824dce2..3e858fa9512 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -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,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);
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index cfa6483c7ae..e3103cea1c3 100644
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -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)
diff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c
b/gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c
new file mode 100644
index 00000000000..29541abd525
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c
@@ -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" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c
b/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c
new file mode 100644
index 00000000000..78a2d786248
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c
@@ -0,0 +1,37 @@
+/* { 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 '\\\)' before numeric constant" } */
+ ; /* { dg-error
"expected '#pragma omp' clause before ':' token" "" { target *-*-* } .-1 } */
+ #pragma omp target uses_allocators (memspace(omp_no_such_space) : foo, bar)
/* { dg-error "modifiers can only be used with a single allocator in
'uses_allocators' clause" } */
+ ;
/* { dg-error "memspace modifier must be constant enum of
'omp_memspace_handle_t' type" "" { target c } .-1 } */
+
/* { dg-error "'omp_no_such_space' has not been declared" "" { target c++ } .-2
} */
+ #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
"expected ':' before ',' token" } */
+ ;
+ #pragma omp target uses_allocators (traitz(traits_array),
memspace(omp_high_bw_mem_space) : bar) /* { dg-error "'traitz' undeclared" "" {
target c } } */
+ ;
/* { dg-error "'memspace' undeclared" "" { target c } .-1 }
*/
+
/* { dg-error "expected '\\\)' before ':' token" "" {
target c } .-2 } */
+
/* { dg-error "'traitz' has not been declared" "" { target
c++ } .-3 } */
+
/* { dg-error "'memspace' has not been declared" "" {
target c++ } .-4 } */
+ return 0;
+}
diff --git a/gcc/testsuite/gfortran.dg/gomp/uses_allocators-1.f90
b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-1.f90
new file mode 100644
index 00000000000..4ca76e7004c
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-1.f90
@@ -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" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/uses_allocators-2.f90
b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-2.f90
new file mode 100644
index 00000000000..530d604902f
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-2.f90
@@ -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
diff --git a/gcc/testsuite/gfortran.dg/gomp/uses_allocators-3.f90
b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-3.f90
new file mode 100644
index 00000000000..064ccf455b1
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-3.f90
@@ -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
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index f1c2b6413a3..7ac0b47ac2d 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -522,6 +522,9 @@ enum omp_clause_code {
/* OpenACC clause: nohost. */
OMP_CLAUSE_NOHOST,
+
+ /* OpenMP clause: uses_allocators. */
+ OMP_CLAUSE_USES_ALLOCATORS,
};
#undef DEFTREESTRUCT
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 99af977979d..a46db024157 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -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(");
{
diff --git a/gcc/tree.cc b/gcc/tree.cc
index 4cf3785270b..973a8366372 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -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
diff --git a/gcc/tree.h b/gcc/tree.h
index 8844471e9a5..bfe2cd82232 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -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)
