| Message ID | b2e9ab91-0996-9f62-64bb-02219a6f7d34@codesourcery.com |
|---|---|
| State | New |
| Headers | show |
| Series | [OpenMP,v2] Implement uses_allocators clause for target regions | expand |
On Tue, May 10, 2022 at 07:29:23PM +0800, Chung-Lin Tang wrote:
> I've attached v2 of the patch. Currently in testing.
Just a general rant, the non-requires dynamic_allocators support
seems to be a total mess in the standard. Probably something
that should be discussed on omp-lang.
Allocators can appear in various places, with requires dynamic_allocators
it is all clear, the only allocators required to be constant expressions
(aka predefined allocators) are on allocate for non-automatic variables
(my understanding is that omp_null_allocator isn't valid for those but I
could be wrong), but there are no other requirements imposed (except of
course referencing a destroyed or not yet initialized allocator is UB),
in particular omp_alloc etc. can be passed omp_null_allocator, or a
variable, and similarly for allocate clause etc.
Without requires dynamic_allocators, there are various extra restrictions
imposed:
1) omp_init_allocator/omp_destroy_allocator may not be called (except for
implicit calls to it from uses_allocators) in a target region
2) omp_alloc etc. can't be called with omp_null_allocator and the argument
has to be a constant expression for a predefined memory allocator
(that is also mentioned on uses_allocators, though that doesn't have to
be visible in the source because it could be lexically included in
a target construct's body)
3) for allocate directive on static vars the above applies plus it has
to be mentioned in uses_allocators
4) for allocate clause e.g. when privatizing stuff or allocate directive
for automatic vars no such restrictions exist
Now, that means that e.g. the user provided uses_allocators without
requires dynamic_allocators are only useful for the case 4), it is unclear
if that was really intended.
With uses_allocators in particular, it is unclear if
uses_allocators(omp_null_allocator) is allowed or not, IMHO it shouldn't,
but I really don't see a restriction disallowing it.
Then there is the issue that 5.0/5.1 said for C/C++ that traits-array
should be
"an identifier of const omp_alloctrait_t * type."
which is wrong for multiple reasons, because identifiers don't have type,
expressions or variables do, but more importantly because from the pointer
to const omp_alloctrait_t it is impossible to find out how many elements
the traits have. 5.2 fixed that to say that it must be an array
(so we thankfully know the size), so we certainly should consider that
change like a defect report against 5.0/5.1 too and require even in the
old syntax an array. Note, I'm afraid we need to support even VLAs,
not just constant size arrays.
There is also in the spec that when allocator in uses_allocators is
a variable, it is treated as a private variable that can't be explicitly
privatized, but nothing said about the traits array, so is say:
void foo () {
omp_allocator_handle_t h;
omp_alloctrait_t t[3] = { ... };
#pragma omp target uses_allocators(h(t)) firstprivate(t)
{
}
ok or not? We need to firstprivatize t so that we can call
h = omp_init_allocator (omp_default_mem_space, 3, t); in the target region
and it is kind of difficult to privatize the same var multiple times.
And yet another issue, in omp_alloctrait_t one can point to other allocators
(with { omp_atk_fallback, some_alloc }). If some_alloc is a predefined
allocator, fine, I don't see big deal with that, especially if that
predefined allocator is also mentioned in uses_allocators clause (before or
after). But if it is a user allocator, there is no restriction on that, and
no way how to map that, say that there should be some specific ordering
of uses_allocators induced omp_init_allocator calls and that we should
somehow replace the host value with privatized target replacement.
More on the actual patch later.
Jakub
On 19/05/2022 17:00, Jakub Jelinek wrote: > Without requires dynamic_allocators, there are various extra restrictions > imposed: > 1) omp_init_allocator/omp_destroy_allocator may not be called (except for > implicit calls to it from uses_allocators) in a target region I interpreted that more like "omp_init_allocator/... is not required to work", as in the set-up steps provided by dynamic_allocators/uses_allocators won't be available. Since we don't have any such on/off mode I don't believe we need to worry about this (and adding extra logic for this is make-work which will not improve the user-experience). > 2) omp_alloc etc. can't be called with omp_null_allocator and the argument > has to be a constant expression for a predefined memory allocator > (that is also mentioned on uses_allocators, though that doesn't have to > be visible in the source because it could be lexically included in > a target construct's body) Again, does a conforming implementation reject this, or is it merely not required to accept it? > 3) for allocate directive on static vars the above applies plus it has > to be mentioned in uses_allocators > 4) for allocate clause e.g. when privatizing stuff or allocate directive > for automatic vars no such restrictions exist > > Now, that means that e.g. the user provided uses_allocators without > requires dynamic_allocators are only useful for the case 4), it is unclear > if that was really intended. > > With uses_allocators in particular, it is unclear if > uses_allocators(omp_null_allocator) is allowed or not, IMHO it shouldn't, > but I really don't see a restriction disallowing it. > > Then there is the issue that 5.0/5.1 said for C/C++ that traits-array > should be > "an identifier of const omp_alloctrait_t * type." > which is wrong for multiple reasons, because identifiers don't have type, > expressions or variables do, but more importantly because from the pointer > to const omp_alloctrait_t it is impossible to find out how many elements > the traits have. 5.2 fixed that to say that it must be an array > (so we thankfully know the size), so we certainly should consider that > change like a defect report against 5.0/5.1 too and require even in the > old syntax an array. Note, I'm afraid we need to support even VLAs, > not just constant size arrays. We are only implementing 5.0 at this time. If 5.2 is less work and the only way to achieve correctness then maybe that's the way to go, but in general, one step at a time, please. > There is also in the spec that when allocator in uses_allocators is > a variable, it is treated as a private variable that can't be explicitly > privatized, but nothing said about the traits array, so is say: > void foo () { > omp_allocator_handle_t h; > omp_alloctrait_t t[3] = { ... }; > #pragma omp target uses_allocators(h(t)) firstprivate(t) > { > } > ok or not? We need to firstprivatize t so that we can call > h = omp_init_allocator (omp_default_mem_space, 3, t); in the target region > and it is kind of difficult to privatize the same var multiple times. > > And yet another issue, in omp_alloctrait_t one can point to other allocators > (with { omp_atk_fallback, some_alloc }). If some_alloc is a predefined > allocator, fine, I don't see big deal with that, especially if that > predefined allocator is also mentioned in uses_allocators clause (before or > after). But if it is a user allocator, there is no restriction on that, and > no way how to map that, say that there should be some specific ordering > of uses_allocators induced omp_init_allocator calls and that we should > somehow replace the host value with privatized target replacement. > > More on the actual patch later. Thank you.
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: > + 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;; 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. > + 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) I think it would be better to have bool variable whether it was memspace or traits modifier, so strcmp just once with each string, not multiple times. In the 5.2 syntax, for memspace it is clear that it has to be an identifier that is the predefined namespace, but for traits it just says it is a variable of alloctrait array type, it is unclear if it must be an identifier or could be say structure element or array element etc. Guess something to discuss on omp-lang and for now pretend it must be an identifier. > + 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; > + } 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. > + 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"); 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. And, I'd strongly prefer to just c_parser_omp_variable_list at this point for the rest if there were modifiers and just fill in OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE and OMP_CLAUSE_USES_ALLOCATORS_TRAITS on each similarly to how e.g. linear or other clause with modifiers are handled. The no modifiers case of course needs its own code so that it handles the ()s. > + 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. > + { > + 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; > + 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. But break; afterwards to avoid the code below. Then, do a VAR_DECL/PARM_DECL check, complain if it is not (see other spots in the function that do check that). Then the bitmap_bit_p stuff above. > + 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[]; 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. 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. Jakub
On Thu, May 19, 2022 at 06:02:43PM +0100, Andrew Stubbs wrote: > On 19/05/2022 17:00, Jakub Jelinek wrote: > > Without requires dynamic_allocators, there are various extra restrictions > > imposed: > > 1) omp_init_allocator/omp_destroy_allocator may not be called (except for > > implicit calls to it from uses_allocators) in a target region > > I interpreted that more like "omp_init_allocator/... is not required to > work", as in the set-up steps provided by dynamic_allocators/uses_allocators > won't be available. Since we don't have any such on/off mode I don't believe > we need to worry about this (and adding extra logic for this is make-work > which will not improve the user-experience). Unfortunately OpenMP as the standard doesn't bother too much with distinctions that e.g. the C/C++ standards make, whether something makes the TU invalid or whether something is only invalid at runtime when reaching it. In any case, I think it would be nice if we diagnosed such uses, doesn't need to be an error, warning would be fine, but help users write portable code, either that they requires dynamic_allocators, or don't and limit themselves to what the standard says should be used in that case. I guess a warning might be better, because we really don't know if it will be actually called at runtime or not from the target region. > > 2) omp_alloc etc. can't be called with omp_null_allocator and the argument > > has to be a constant expression for a predefined memory allocator > > (that is also mentioned on uses_allocators, though that doesn't have to > > be visible in the source because it could be lexically included in > > a target construct's body) > > Again, does a conforming implementation reject this, or is it merely not > required to accept it? I think it is another fuzzy area, but again it would be nice to at least get warnings. Jakub
Hi Jakub, On 19.05.22 18:00, Jakub Jelinek wrote: > On Tue, May 10, 2022 at 07:29:23PM +0800, Chung-Lin Tang wrote: >> I've attached v2 of the patch. Currently in testing. > Just a general rant, the non-requires dynamic_allocators support > seems to be a total mess in the standard. Probably something > that should be discussed on omp-lang. Or in some issue. Some newer developments (all links unfortunately nonpublic): There is now a nearly ready example for the 5.2.1 example document, cf. https://github.com/OpenMP/examples-internal/issues/275 https://github.com/OpenMP/spec/issues/3229 improves the wording related to 'requires dynamic_allocators' – vote was after the 5.2 release. > [...]Allocators can appear in various places, with requires dynamic_allocators > it is all clear, the only allocators required to be constant expressions > (aka predefined allocators) are on allocate for non-automatic variables Side remark: the OMP_ALLOCATOR environment variable sets the def-allocator-var ICV – but besides pre-defined allocators, it also permits to define (traits, memspace, ...) a new default allocator (new in OMP 5.1). And this one can then seemingly also be used in the target region (only with 'requires dynamic_allocators'). I note that GCC supports OMP_ALLOCATOR but seemingly only with predefined allocators (→ OMP 5.0). – I think we need to open PR for that one and/or a new line in the 5.1 implementation tables. > (my understanding is that omp_null_allocator isn't valid for those but I > could be wrong), (I read it likewise as it is not predefined,) > Without requires dynamic_allocators, there are various extra restrictions > imposed: > 1) omp_init_allocator/omp_destroy_allocator [...] > 2) omp_alloc etc. [...] > 3) for allocate directive on static vars [...] > 4) for allocate clause e.g. when privatizing stuff or allocate directive > for automatic vars no such restrictions exist > > Now, that means that e.g. the user provided uses_allocators without > requires dynamic_allocators are only useful for the case 4), it is unclear > if that was really intended. Note that (4) not only applies for 'allocate' on the 'target' construct but it can be also be used on any other directive inside the target construct, i.e.: #pragma omp target uses_allocators(omp_cgroup_mem_alloc) #pragma omp teams reduction(+:xbuf) thread_limit(N) \ allocate(omp_cgroup_mem_alloc:xbuf) (from the example issue, linked above). > With uses_allocators in particular, it is unclear if > uses_allocators(omp_null_allocator) is allowed or not, IMHO it shouldn't, > but I really don't see a restriction disallowing it. I think it does not count as predefined allocator and the (new, post-5.2) wording makes clear that the default allocator (which is associated with the omp_null_allocator per wording in, e.g., omp_alloc) is only valid with 'dynamic_allocators'. > Then there is the issue that 5.0/5.1 said for C/C++ that traits-array > should be > "an identifier of const omp_alloctrait_t * type." > which is wrong for multiple reasons, because identifiers don't have type, > expressions or variables do, but more importantly because from the pointer > to const omp_alloctrait_t it is impossible to find out how many elements > the traits have. 5.2 fixed that to say that it must be an array > (so we thankfully know the size), so we certainly should consider that > change like a defect report against 5.0/5.1 too and require even in the > old syntax an array. Note, I'm afraid we need to support even VLAs, > not just constant size arrays. I concur that 5.2 should be regarded both as fix of old bugs and syntax extension. (Additionally, as we already support the 5.2 syntax.) There were some improvements discussed/proposed in https://github.com/OpenMP/spec/issues/3285 It is a bit difficult to read as I confused two things at the beginning (mixing allocator expression / traits argument when reading the bullet points). – But it relates to some of the items you raised here. > There is also in the spec that when allocator in uses_allocators is > a variable, it is treated as a private variable that can't be explicitly > privatized, but nothing said about the traits array, so is say: > void foo () { > omp_allocator_handle_t h; > omp_alloctrait_t t[3] = { ... }; > #pragma omp target uses_allocators(h(t)) firstprivate(t) > { > } > ok or not? (try in addition 'allocate(h : t)' ) > We need to firstprivatize t so that we can call > h = omp_init_allocator (omp_default_mem_space, 3, t); in the target region > and it is kind of difficult to privatize the same var multiple times. I think this relates to the generic question related to mapping/firstprivatizing const/constexpr/PARAMETER etc. https://github.com/OpenMP/spec/issues/2158 which really should be fixed. > And yet another issue, in omp_alloctrait_t one can point to other allocators > (with { omp_atk_fallback, some_alloc }). If some_alloc is a predefined > allocator, fine, I don't see big deal with that, especially if that > predefined allocator is also mentioned in uses_allocators clause (before or > after). But if it is a user allocator, there is no restriction on that, and > no way how to map that, say that there should be some specific ordering > of uses_allocators induced omp_init_allocator calls and that we should > somehow replace the host value with privatized target replacement. I think that might need a clarification/fix on the OpenMP spec side. Tobias ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
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. 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)
On Mon, May 30, 2022 at 10:43:30PM +0800, Chung-Lin Tang wrote: > > 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. I guess this should be raised on omp-lang then what we really want. Because the 5.1 syntax definitely allowed multiple allocators. > > 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. Sure, but the spec says that omp_alloc (42, omp_null_allocator) is invalid in target regions unless requires dynamic_allocators is seen first. And uses_allocators (omp_null_allocator) shouldn't make that valid. > @@ -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); Each vec/auto_vec with a new type brings quite some overhead, a lot of functions need to be instantiated for it. I think it would be far easier to use raw token parsing: unsigned int pos = 1; bool has_modifiers = false; while (true) { c_token *tok = c_parser_peek_nth_token_raw (parser, pos); if (tok->type != CPP_NAME) break; ++pos; if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_OPEN_PAREN) { ++pos; if (!c_parser_check_balanced_raw_token_sequence (parser, &pos) || c_parser_peek_nth_token_raw (parser, pos)->type != CPP_CLOSE_PAREN) break; ++pos; } tok = c_parser_peek_nth_token_raw (parser, pos); if (tok->type == CPP_COLON) { has_modifiers = true; break; } if (tok->type != CPP_COMMA) break; ++pos; } This should (haven't tested it though, so sorry if there are errors) cheaply determine if one should or shouldn't parse modifiers and then can just do parsing of modifiers if (has_modifiers) and then just the list (with ()s inside of list only if (!has_modifiers)). > @@ -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)\ Can you please fix up both the IS_DEVICE_PTR and newly added HAS_DEVICE_ADDR change to have a space before \ ? > + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USES_ALLOCATORS)) > +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); In C++ FE one can use tentative parsing instead, but nth_token will work too. size_t pos = 1; bool has_modifiers = false; while (true) { if (!cp_lexer_nth_token_is (parser, pos, CPP_NAME)) break; ++pos; if (cp_lexer_nth_token_is (parser, pos, CPP_OPEN_PAREN)) { size_t npos = cp_parser_skip_balanced_tokens (parser, pos); if (npos == pos) break; pos = npos; } cp_token *tok = cp_lexer_peek_nth_token (parser, pos); if (tok->type == CPP_COLON) { has_modifiers = true; break; } if (tok->type != CPP_COMMA) break; ++pos; } > @@ -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)\ Ditto. > + 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); Are you sure about the above line? Should we allow omp_allocator_handle_t & type vars in the list? > + if (TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE > + || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE (t))), > + "omp_allocator_handle_t") != 0) On the other side, if type_dependent_expression_p (t), I think we shouldn't diagnose this but postpone it till instantiation. And there should be in the testsuite a C++ testcase, where it uses_allocators inside of a template, in one place with a non-dependent omp_allocator_handle_t type, in another case e.g. with template parameter type that is later instantiated with omp_allocator_handle_t type. > + { > + 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; As I said earlier, I'd prefer to keep them and if for now you don't want to warn for uses of allocators that aren't mentioned in uses_allocators, just ignore them when actually privatizing them. > + > + 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"); But of course this case should have remove = true; Jakub
On Mon, May 30, 2022 at 07:23:55PM +0200, Jakub Jelinek via Gcc-patches wrote: > On Mon, May 30, 2022 at 10:43:30PM +0800, Chung-Lin Tang wrote: > > > 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. > > I guess this should be raised on omp-lang then what we really want. > Because the 5.1 syntax definitely allowed multiple allocators. The response I got on omp-lang is that it is intentional that in the new syntax only a single allocator is allowed. So I'd suggest to implement: 1) if has_modifiers (i.e. certainly new syntax), only allow a single enumerator / identifier for a variable and no ()s after it 2) if !has_modifiers and there is exactly one allocator without ()s, treat it like new syntax 3) otherwise, it is the old (5.1) syntax, which allows a list and that list can contain ()s for traits, but in the light of the 5.2 wording, I'd even for that case avoid diagnosing missing traits for non-predefined allocators 4) omp_null_allocator should be diagnosed as invalid, private (omp_null_allocator) is rejected... 5) for C++, we should handle FIELD_DECLs, but it shouldn't be hard, just look how it is handled for private too Jakub
On 2022/5/31 6:02 PM, Jakub Jelinek wrote: > 5) for C++, we should handle FIELD_DECLs, but it shouldn't be hard, just > look how it is handled for private too > > Jakub About private() for non-static members, is it really working right now? A simple test: struct C { omp_allocator_handle_t a; void foo (void) { #pragma omp target private (a) a = (omp_allocator_handle_t) 0; } }; int main (void) { C c; c.foo (); return 0; } After C++ front-end processing we get: { omp_allocator_handle_t D.2823 [value-expr: ((struct C *) this)->a]; #pragma omp target private(D.2823) { { <<cleanup_point <<< Unknown tree: expr_stmt (void) (D.2823 = 0) >>>>>; } } } The OMP field privatization seems to be doing something here. However gimplify turns this into: void C::foo (struct C * const this) { omp_allocator_handle_t a [value-expr: ((struct C *) this)->a]; #pragma omp target num_teams(1) thread_limit(0) private(a) \ map(alloc:MEM[(char *)this] [len: 0]) map(firstprivate:this [pointer assign, bias: 0]) { this->a = 0; } } This doesn't look quite right for private clause. I don't quite expect a zero-length mapping of this[:0], nor reverting the gimple to use "this->a" for a private copy. Chung-Lin
On Mon, Jun 06, 2022 at 09:19:18PM +0800, Chung-Lin Tang wrote: > On 2022/5/31 6:02 PM, Jakub Jelinek wrote: > > 5) for C++, we should handle FIELD_DECLs, but it shouldn't be hard, just > > look how it is handled for private too > > > > Jakub > > About private() for non-static members, is it really working right now? Perhaps we have a bug that we should file in bugzilla and should fix. Can you try omp parallel or omp target in the test instead? > A simple test: > > struct C { > omp_allocator_handle_t a; > void foo (void) { > #pragma omp target private (a) > a = (omp_allocator_handle_t) 0; > } > }; > > int main (void) > { > C c; > c.foo (); > return 0; > } Jakub
On 2022/6/6 9:22 下午, Jakub Jelinek wrote: > On Mon, Jun 06, 2022 at 09:19:18PM +0800, Chung-Lin Tang wrote: >> On 2022/5/31 6:02 PM, Jakub Jelinek wrote: >>> 5) for C++, we should handle FIELD_DECLs, but it shouldn't be hard, just >>> look how it is handled for private too >>> >>> Jakub >> >> About private() for non-static members, is it really working right now? > > Perhaps we have a bug that we should file in bugzilla and should fix. > > Can you try omp parallel or omp target in the test instead? I see it works for omp parallel/task, gimplify results: void C::foo (struct C * const this) { omp_allocator_handle_t a [value-expr: ((struct C *) this)->a]; #pragma omp parallel private(a) { a = 0; } } I'll file a bugzilla for the target construct. That said, can we delay FIELD_DECL support for uses_allocators? (which is target construct only) Since it appears to be not trivial at the moment. Thanks, Chung-Lin >> A simple test: >> >> struct C { >> omp_allocator_handle_t a; >> void foo (void) { >> #pragma omp target private (a) >> a = (omp_allocator_handle_t) 0; >> } >> }; >> >> int main (void) >> { >> C c; >> c.foo (); >> return 0; >> } > > Jakub >
On Mon, Jun 06, 2022 at 09:38:30PM +0800, Chung-Lin Tang wrote: > I'll file a bugzilla for the target construct. Thanks. > That said, can we delay FIELD_DECL support for uses_allocators? (which is target construct only) > Since it appears to be not trivial at the moment. Sure. But would be nice to file a PR to track it once the patch is committed. Jakub
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)