Re: [PATCH, OpenMP, v2] Implement uses_allocators clause for target regions

[Tobias Burnus <tobias@codesourcery.com>]

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