From: Chung-Lin Tang <cltang@codesourcery.com>
To: Jakub Jelinek <jakub@redhat.com>
Cc: Tobias Burnus <tobias@codesourcery.com>,
gcc-patches <gcc-patches@gcc.gnu.org>,
Catherine Moore <clm@codesourcery.com>,
Andrew Stubbs <ams@codesourcery.com>,
Hafiz Abid Qadeer <abid_qadeer@mentor.com>
Subject: Re: [PATCH, OpenMP, v2] Implement uses_allocators clause for target regions
Date: Mon, 30 May 2022 22:43:30 +0800 [thread overview]
Message-ID: <f43b4e82-4f21-ca3b-4f5c-e7edc942d231@codesourcery.com> (raw)
In-Reply-To: <YoaCj4ZNKF7Opqem@tucnak>
[-- Attachment #1: Type: text/plain, Size: 11560 bytes --]
Hi Jakub,
this is v3 of the uses_allocators patch.
On 2022/5/20 1:46 AM, Jakub Jelinek wrote:
> On Tue, May 10, 2022 at 07:29:23PM +0800, Chung-Lin Tang wrote:
>> @@ -15624,6 +15626,233 @@ c_parser_omp_clause_allocate (c_parser *parser, tree list)
>> return nl;
>> }
>>
>> +/* OpenMP 5.2:
>> + uses_allocators ( allocator-list )
>
> As uses_allocators is a 5.0 feature already, the above should say
> /* OpenMP 5.0:
>> +
>> + allocator-list:
>> + allocator
>> + allocator , allocator-list
>> + allocator ( traits-array )
>> + allocator ( traits-array ) , allocator-list
>> +
>
> And here it should add
> OpenMP 5.2:
Done.
>> + 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;;
>
> Double ;;, should be just ;
> But more importantly, it is more complex.
> When you see
> uses_allocators(traits or
> uses_allocators(memspace
> it is not given that it has modifiers. While the 5.0/5.1 syntax had
> a restriction that when allocator is not a predefined allocator (and
> traits or memspace aren't predefined allocators) it must use ()s with
> traits, so
> uses_allocators(traits)
> uses_allocators(memspace)
> uses_allocators(traits,memspace)
> are all invalid,
> omp_allocator_handle_t traits;
> uses_allocators(traits(mytraits))
> or
> omp_allocator_handle_t memspace;
> uses_allocators(memspace(mytraits),omp_default_mem_alloc)
> are valid in the old syntax.
>
> So, I'm afraid to find out if the traits or memspace identifier
> seen after uses_allocator ( are modifiers or not we need to
> peek (for C with c_parser_peek_nth_token_raw) through all the
> modifiers whether we see a : and only in that case say they
> are modifiers rather than the old style syntax.
The parser parts have been rewritten to allow this kind of use now.
New code essentially parses lists of "id(id), id(id), ...", possibly delimited
by a ':' marking the modifier/allocator lists.
> I don't really like the modifiers handling not done in a loop.
> As I said above, there needs to be some check whether there are modifiers or
> not, but once we figure out there are modifiers, it should be done in a loop
> with say some mask var on which traits have been already handled to diagnose
> duplicates, we don't want to do the parsing code twice.
Now everything is done in loops. The new code should be considerably simpler now.
> This feels like you only accept a single allocator in the new syntax,
> but that isn't my reading of the spec, I'd understand it as:
> uses_allocators (memspace(omp_high_bw_mem_space), traits(foo_traits) : bar, baz, qux)
> being valid too.
This patch now allows multiple allocators to be specified in new syntax, although I have
to note that the 5.2 specification of uses_allocators (page 181) specifically says
"allocator: expression of allocator_handle_type" for the "Arguments" description,
not a "list" like the allocate clause.
>> + 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)))
>
> You can't just use DECL_UID before you actually verify it is a variable.
> So IMHO this particular if should be moved down somewhat.
Guarded now.
>> + {
>> + 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;
>> + }
>
> I'd add break; after remove = true;
Added some such breaks.
>> + 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;
>> + }
>
> It should be only removed if we emit the error (again with break; too).
> IMHO (see the other mail) we should complain here if it has value 0
> (the omp_null_allocator case), dunno if we should error or just warn
> if the value is outside of the range of known predefined identifiers (1..8
> currently I think). > But, otherwise, IMHO we need to keep it around, perhaps replacing the
> CONST_DECL with INTEGER_CST, for the purposes of checking what predefined
> allocators are used in the region.
omp_alloc in libgomp does handle the omp_null_allocator case, by converting it
to something else.
The code already checks if the type is the OpenMP specified 'omp_memspace_handle_t'
enumeration type. A CONST_DECL of that type should be guaranteed a pre-defined
identifier.
>> + 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;
>> + }
>
> Again, wonder if it shouldn't after checking it replace the CONST_DECL with
> an INTEGER_CST for the purposes of the middle-end.
>
>> + 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))
>
> I'd diagnose if the array is incomplete, say
> extern omp_alloctrait_t traits[];
I've added a DECL_SIZE == NULL_TREE check, although the TREE_READONLY check for
the element type usually seems to already disqualify this case (because of the 'extern')
> For the 5.2 syntax, there is also the restriction that
> "be defined in the same scope as the construct on which the clause appears"
> which I don't see being checked here. Unclear whether it applies to the
> old syntax too.
>
> But again, it should also check that it is a VAR_DECL, it isn't extern
> etc.
Our interpretation of the requirement of "same-scope", and that it must be a constant array,
is that the traits array is intended to be inlined into the target region (instead of
more hassle issues related to transporting it to the offload target), which is what we
do right now.
In this case "same scope" is probably a little bit of overkill, it only needs to be staticly
known/computable by the compiler.
> For C++, I have to wonder if at allocator couldn't be a non-static data
> member of a class inside of methods, that is something that can be generally
> privatized too.
Maybe later, I've sorry'ed FIELD_DECLs for now.
Asides from the review issues, this patch also includes some fixes for the allocate clause
firstprivate transfering on task constructs, triggered by the changes in the Fortran FE.
Tested without regressions on mainline.
Thanks,
Chung-Lin
2022-05-30 Chung-Lin Tang <cltang@codesourcery.com>
gcc/c-family/ChangeLog:
* c-omp.cc (c_omp_split_clauses): Add OMP_CLAUSE_USES_ALLOCATORS case.
* c-pragma.h (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_USES_ALLOCATORS.
gcc/c/ChangeLog:
* c-parser.cc (c_parser_omp_clause_name): Add case for uses_allocators
clause.
(c_parser_omp_clause_uses_allocators): New function.
(c_parser_omp_all_clauses): Add PRAGMA_OMP_CLAUSE_USES_ALLOCATORS case.
(OMP_TARGET_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USES_ALLOCATORS to mask.
* c-typeck.cc (c_finish_omp_clauses): Add case handling for
OMP_CLAUSE_USES_ALLOCATORS.
gcc/cp/ChangeLog:
* parser.cc (cp_parser_omp_clause_name): Add case for uses_allocators
clause.
(cp_parser_omp_clause_uses_allocators): New function.
(cp_parser_omp_all_clauses): Add PRAGMA_OMP_CLAUSE_USES_ALLOCATORS case.
(OMP_TARGET_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USES_ALLOCATORS to mask.
* semantics.cc (finish_omp_clauses): Add case handling for
OMP_CLAUSE_USES_ALLOCATORS.
fortran/ChangeLog:
* gfortran.h (struct gfc_omp_namelist): Add memspace_sym, traits_sym
fields.
(OMP_LIST_USES_ALLOCATORS): New list enum.
* openmp.cc (enum omp_mask2): Add OMP_CLAUSE_USES_ALLOCATORS.
(gfc_match_omp_clause_uses_allocators): New function.
(gfc_match_omp_clauses): Add case to handle OMP_CLAUSE_USES_ALLOCATORS.
(OMP_TARGET_CLAUSES): Add OMP_CLAUSE_USES_ALLOCATORS.
(resolve_omp_clauses): Add "USES_ALLOCATORS" to clause_names[].
* dump-parse-tree.cc (show_omp_namelist): Handle OMP_LIST_USES_ALLOCATORS.
(show_omp_clauses): Likewise.
* trans-array.cc (gfc_conv_array_initializer): Adjust array index
to always be a created tree expression instead of NULL_TREE when zero.
* trans-openmp.cc (gfc_trans_omp_clauses): For ALLOCATE clause, handle
using gfc_trans_omp_variable for EXPR_VARIABLE exprs.
Add handling of OMP_LIST_USES_ALLOCATORS case.
* types.def (BT_FN_VOID_PTRMODE): Define.
(BT_FN_PTRMODE_PTRMODE_INT_PTR): Define.
gcc/ChangeLog:
* builtin-types.def (BT_FN_VOID_PTRMODE): Define.
(BT_FN_PTRMODE_PTRMODE_INT_PTR): Define.
* omp-builtins.def (BUILT_IN_OMP_INIT_ALLOCATOR): Define.
(BUILT_IN_OMP_DESTROY_ALLOCATOR): Define.
* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_USES_ALLOCATORS.
* tree-pretty-print.cc (dump_omp_clause): Handle OMP_CLAUSE_USES_ALLOCATORS.
* tree.h (OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR): New macro.
(OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE): New macro.
(OMP_CLAUSE_USES_ALLOCATORS_TRAITS): New macro.
* tree.cc (omp_clause_num_ops): Add OMP_CLAUSE_USES_ALLOCATORS.
(omp_clause_code_name): Add "uses_allocators".
* gimplify.cc (gimplify_scan_omp_clauses): Add checking of OpenMP target
region allocate clauses, to require a uses_allocators clause to exist
for allocators.
(gimplify_omp_workshare): Add handling of OMP_CLAUSE_USES_ALLOCATORS
for OpenMP target regions; create calls of omp_init/destroy_allocator
around target region body.
* omp-low.cc (lower_private_allocate): Adjust receiving of allocator.
(lower_rec_input_clauses): Likewise.
(create_task_copyfn): Add dereference for allocator if needed.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/uses_allocators-1.c: New test.
* c-c++-common/gomp/uses_allocators-2.c: New test.
* c-c++-common/gomp/uses_allocators-3.c: New test.
* gfortran.dg/gomp/allocate-1.f90: Adjust testcase.
* gfortran.dg/gomp/uses_allocators-1.f90: New test.
* gfortran.dg/gomp/uses_allocators-2.f90: New test.
* gfortran.dg/gomp/uses_allocators-3.f90: New test.
[-- Attachment #2: uses_allocators-v3.patch --]
[-- Type: text/plain, Size: 66774 bytes --]
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 66d17a2673d..50db6936728 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -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:
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 492d995a281..2022c16802d 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -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,199 @@ 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))
+ {
+ if (modifiers)
+ {
+ c_parser_error (parser, "modifiers cannot be used with "
+ "(deprecated) traits array list syntax");
+ goto end;
+ }
+ 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)
+ 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 +17274,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 +21292,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 4f3611f1b89..bcd4ca7074b 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -14763,6 +14763,110 @@ 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 (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)
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 868b8610d60..8a0b7783e1c 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -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,220 @@ 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))
+ {
+ if (modifiers)
+ {
+ cp_parser_error (parser, "modifiers cannot be used with "
+ "(deprecated) traits array list syntax");
+ goto end;
+ }
+ 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)
+ 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 +40669,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 +44684,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 cd7a2818feb..eba786ae5af 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -7766,6 +7766,104 @@ 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 (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)
diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc
index 4e8986bd599..3f1396544ca 100644
--- a/gcc/fortran/dump-parse-tree.cc
+++ b/gcc/fortran/dump-parse-tree.cc
@@ -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 ();
}
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 5d970bc1df0..b02cbf87048 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -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. */
};
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 63fd4dd2767..df4e5d0588c 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -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,500 @@ 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)
+ {
+ if (modifiers)
+ {
+ gfc_error ("Modifiers cannot be used with (deprecated) traits "
+ "array list syntax at %C");
+ goto error;
+ }
+ 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)
+ for (unsigned i = 0; i < allocators->length (); i++)
+ {
+ item& it = (*allocators)[i];
+
+ gfc_symbol *allocator_sym;
+
+ 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;
+ }
+
+ 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 %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 = 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 +3466,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 +4195,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 +6828,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..47c636420d3 100644
--- a/gcc/fortran/trans-array.cc
+++ b/gcc/fortran/trans-array.cc
@@ -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)
{
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index bfd24f964ea..1d501d6f720 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -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];
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 cd1796643d7..16a495ba586 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -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);
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index ee5213eedcf..ccfd6a542ec 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/omp-low.cc b/gcc/omp-low.cc
index 16f596587e8..e2784a24232 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -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);
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..e57b2906c83
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c
@@ -0,0 +1,33 @@
+/* { 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 "memspace modifier must be constant enum of 'omp_memspace_handle_t' type" "" { target c } } */
+ ; /* { 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'" } */
+ ;
+ return 0;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-3.c b/gcc/testsuite/c-c++-common/gomp/uses_allocators-3.c
new file mode 100644
index 00000000000..80b2844729a
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-3.c
@@ -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" } } */
diff --git a/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90 b/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90
index 8bc6b768778..f5707899eff 100644
--- a/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90
@@ -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);
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..3762250ee44
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-2.f90
@@ -0,0 +1,40 @@
+! { 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
+
+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..0f024264700
--- /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'' requires 'uses_allocators.bar.' clause in target region" }
+ block
+ allocate(arr(100))
+ end block
+
+end program main
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 2383b570f49..91300d42b5a 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 333ac23aeb2..5f5a15b48db 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 df441c6b223..407233a5b6d 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 c92c5bf344b..c31cc9ce859 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -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)
next prev parent reply other threads:[~2022-05-30 14:43 UTC|newest]
Thread overview: 20+ messages / expand[flat|nested] mbox.gz Atom feed top
2022-05-06 13:20 [PATCH, OpenMP] " Chung-Lin Tang
2022-05-06 16:40 ` Tobias Burnus
2022-05-10 11:29 ` [PATCH, OpenMP, v2] " Chung-Lin Tang
2022-05-19 16:00 ` Jakub Jelinek
2022-05-19 17:02 ` Andrew Stubbs
2022-05-19 17:55 ` Jakub Jelinek
2022-05-20 6:59 ` Tobias Burnus
2022-05-19 17:46 ` Jakub Jelinek
2022-05-30 14:43 ` Chung-Lin Tang [this message]
2022-05-30 17:23 ` Jakub Jelinek
2022-05-31 10:02 ` Jakub Jelinek
2022-06-06 13:19 ` Chung-Lin Tang
2022-06-06 13:22 ` Jakub Jelinek
2022-06-06 13:38 ` Chung-Lin Tang
2022-06-06 13:42 ` Jakub Jelinek
2022-06-09 6:21 ` [PATCH, OpenMP, v4] " Chung-Lin Tang
2022-06-09 12:22 ` Jakub Jelinek
2022-06-13 13:29 ` Chung-Lin Tang
2022-06-13 14:04 ` Jakub Jelinek
2023-02-09 11:26 ` [og12] '{c-c++-common,gfortran.dg}/gomp/uses_allocators-*' -> 'libgomp.{c-c++-common,fortran}/uses_allocators-*' (was: [PATCH, OpenMP] Implement uses_allocators clause for target regions) Thomas Schwinge
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=f43b4e82-4f21-ca3b-4f5c-e7edc942d231@codesourcery.com \
--to=cltang@codesourcery.com \
--cc=abid_qadeer@mentor.com \
--cc=ams@codesourcery.com \
--cc=clm@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=jakub@redhat.com \
--cc=tobias@codesourcery.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).