Hi Thomas,
Thomas Schwinge wrotw:
On 2026-09-26T11:01:01+0200, Tobias Burnus<[email protected]> wrote:
Tobias Burnus wrote:
Similar to CUDA language's __SHARED__ attribute that turns static variables
(Actually, not only 'static' ones; it causes additional semantic
changes... But that's not relevant to your case, here.)
(For non-static cases, OpenMP has dyn_groupprivate and many API
routines, cf. below. - But OpenMP's groupprivate is only for static
variable - global or local.)
Hey, in a message on Friday night you said that I have until today to
review this. ;-O
But had I committed it this morning, your email would have also arrived many
hours afterwards, given that it is already 11am ...
--- a/gcc/config/nvptx/nvptx.cc
+++ b/gcc/config/nvptx/nvptx.cc
+ || lookup_attribute ("omp groupprivate", DECL_ATTRIBUTES (decl)))
{
area = DATA_AREA_SHARED;
Yes, sure.
So, a few weeks ago, I was doing a "very similar" thing in a different
context, and I first tried the very same thing that you did here, but
then didn't see the attribure propagated from host to nvptx offloading
Here it works – at least my debug 'warning_at' was called and also the
added testcase fails when 'omp groupprivate' is commented out.
what I then did, was to mark up the variable host-side with attribute
'oacc gang-private' (à la 'nvptx_goacc_adjust_private_decl' as called via
'TARGET_GOACC_ADJUST_PRIVATE_DECL'/'targetm.goacc.adjust_private_decl'
For the final implementation (i.e. the one that adds LDS address space),
I was thinking of using a hook - but I wasn't sure whether
targetm.goacc.adjust_private_decl
was the right one or not in terms of the exact semantic or whether a new
hook is need. As first quick solution, I voted for the committed one.
I currently do
not see us supporting groupprivate on the host for static variables.
If I remember correctly, we had or maybe still have some code in
'libgomp/target.c', to implement 'firstprivate' etc. for OpenMP 'target'
in host-fallback execution (using 'malloc') -- maybe something
(conceptually) similar is applicale for 'groupprivate'?
I think we could implement it as follow. Note: This examples assumes
that the device supports it and the host not. Otherwise, the more generic
solution would be:
#if ACCELERATOR
if (!hook)
#else
if (!initial_device () || !hook)
#endif
In any case, assume a user code like:
int global;
#pragma omp groupprivate(global)
int f(bool first) {
static int local;
#pragma omp groupprivate(local)
if (first)
local = 5;
return ++local;
}
int g() {
return ++global;
}
int foo () {
...
global++;
#pragma omp (target) teams num_teams(n)
{
int x = global;
}
...
It could then be implemented as:
int global; // needs to be marked as 'shared' / LDS on the device side
int g() {
int *_global = omp_is_initial_device()
? __omp_get_static_var (&global, sizeof (int),
omp_cgroup_mem_alloc, alignof (int));
: &global;
...
return ++(*_global);
}
int f(bool first) {
static int local;
int *_local = omp_is_initial_device()
? __omp_get_static_var (&local, sizeof (int),
omp_cgroup_mem_alloc, alignof (int));
: &local;
if (first)
*_local = 5;
return ++(*_local);
}
int foo () {
...
global++; /* direct use if there is 'omp teams' is in this function or
if in 'main'; otherwise like in 'g'. */
#pragma omp (target) teams num_teams(n) // note: groupprivate is never mapped
{
int *_global = omp_is_initial_device()
? __omp_get_static_var (&global, sizeof (int),
omp_cgroup_mem_alloc, alignof
(int));
: &global;
...
int x = *_global;
}
...
For better performance, the current code could be kept for device_type(nohost)
(assumption: all nonhosts support 'groupprivate' in hardware).
For device_type(any), I hope the code above is fast enough - esp. if the pointer
assignment is done to a function local stack variable as '_local = &local'.
Regarding __omp_get_static_var: If not nested in a teams, it would return the
first argument - otherwise, it will allocate memory for all teams - record it in
a by-teams variable in libgomp - and return mem[thread_num()]. If allocated,
it can returned that saved memory - and when 'teams' ends, all memory can be
freed.
Note: 'omp teams' can neither be nested nor run asynchronously.
'omp target teams' currently runs on the host with a single team, but this seems
to be invalid with 'num_teams([lower:]upper)' esp. as lower := upper if not
specified.
* * *
Note: There is also 'dyn_groupprivate' which is related but quite
differently handled (spec wise and implementation wise).
ACK.
For completeness: 'dyn_groupprivate(size)' tells how much memory should be
allocated (per team) and the full usage is then:
#pragma omp target teams dyn_groupprivate(size)
...
int *my_team_local_mem = omp_get_dyn_gprivate_ptr();
There is a fallback-handling if the size is too large both in the 'clause' and
how
the API routine should react. There are also some query functions and there are
actually two API routine, one just returning the value (which might be NULL) the
other allocating memory in this case - if not already done via the launch
function.
In the CUDA language, __SHARED__ is also used for that purpose:
extern __SHARED__ int var[];
the size is specified likewise at the kernel launch:
myKernel<< ..., size>>();
In any case, this needs quite a bit work in libgomp and some builtins
to implement this efficiently (= mostly backend work). On the parser side,
it is implemented (such that a nice 'sorry' is printed in gimplify.cc).
In any case, as it requires some larger work, it will take a while ...
For 'omp target (teams)' a couple of new clauses need to be passed on,
dyn_groupprivate' is only one of them (strict modifier, spatial dimensions, ...)
for which a new ABI is needed.
Tobias
