Hi!
On 2023-11-17T16:24:46+0100, I wrote:
> [...] attached "Add 'libgomp.c++/static-local-variable-1.C'" [...]
Now, working on translating this into an OpenMP 'target' variant. My
goal here is not necessarily to make this work now, but rather to figure
out whether '-fthreadsafe-statics' actually does or doesn't need to be
supported in offloading compilation, whether '__cxa_guard_acquire' is in
fact unreachable there. (Currently the latter symbol isn't available in
offloading compilation; as you know I'm currently working on GPU
libstdc++ library support.) However, GCC offloading compilation
currently fails differently, as follows:
r.cc:16:12: error: variable ‘_ZGVZL1fvE1s’ has been referenced in offloaded
code but hasn’t been marked to be included in the offloaded code
16 | static S s;
| ^
lto1: fatal error: errors during merging of translation units
compilation terminated.
nvptx mkoffload: fatal error:
build-gcc/gcc/x86_64-pc-linux-gnu-accel-nvptx-none-gcc returned 1 exit status
[...]
... with:
$ c++filt _ZGVZL1fvE1s
guard variable for f()::s
Now I wonder how is that supposed to behave; is this valid OpenMP
'target' code at all? Can you please help me find my way through the
OpenMP specification regarding this?
In OpenMP 5.2, 5.1.1 "Variables Referenced in a Construct", we have:
- Variables with static storage duration that are declared in a scope
inside the construct are shared.
Does this apply to a 'declare target'ed function 'f'? (I was thinking:
"dynamic extend" of the scope of the 'target' construct?) Ah, probably
that's 5.1.2 "Variables Referenced in a Region but not in a Construct":
- Variables with static storage duration that are declared in called
routines in the region are shared.
In 7.8 "Declare Target Directives", we have:
If a variable with static storage duration is declared in a device routine
then the named variable is
treated as if it had appeared in an 'enter' clause on a declare target
directive.
Similarly, in 13.8 "'target' Construct":
If a variable with static storage duration is declared in a 'target'
construct that does not specify a
'device' clause in which the 'ancestor' _device-modifier_ appears then the
named variable is
treated as if it had appeared in a 'enter' clause on a declare target
directive.
Per those occurrences, and per GCC not raising an error when encountering
a static local variable, I assume this is intended to work "as expected"?
On the other hand, NVHPC nvc++ 23.1-0 fails, too:
NVC++-S-1062-Support procedure called within a compute region -
__cxa_guard_acquire (r.cc: 16)
[local to r_cc]::f():
16, Accelerator restriction: unsupported call to support routine
'__cxa_guard_acquire'
NVC++/x86-64 Linux 23.1-0: compilation completed with severe errors
Hmm...
Grüße
Thomas
-----------------
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
#pragma omp declare target
struct S
{
S()
{
}
~S()
{
}
};
static void f()
{
static S s;
}
#pragma omp end declare target
int main()
{
#pragma omp target
{
f();
}
}
