https://gcc.gnu.org/bugzilla/show_bug.cgi?id=90779
Jakub Jelinek <jakub at gcc dot gnu.org> changed:
What |Removed |Added
----------------------------------------------------------------------------
CC| |amonakov at gcc dot gnu.org,
| |vries at gcc dot gnu.org
--- Comment #7 from Jakub Jelinek <jakub at gcc dot gnu.org> ---
Something like:
extern int foo (void);
static int var = 5;
#pragma omp declare target to (var)
int
main ()
{
int r = 19;
#pragma omp target map(tofrom: r)
{
r += ++var;
}
#pragma omp target map(tofrom: r)
{
r += ++var;
}
if (r != 19 + 6 + 7)
__builtin_abort ();
foo ();
return 0;
}
in one source and:
static int var = 5;
#pragma omp declare target to (var)
int
foo ()
{
int r = 19;
#pragma omp target map(tofrom: r)
{
r += ++var;
}
#pragma omp target map(tofrom: r)
{
r += ++var;
}
if (r != 19 + 6 + 7)
__builtin_abort ();
return 0;
}
needs to work. Looking at the PTX assembly for this, I see
// BEGIN VAR DEF: var
.global .align 4 .u32 var[1] = { 5 };
//:FUNC_MAP "main$_omp_fn$1"
//:FUNC_MAP "main$_omp_fn$0"
//:VAR_MAP "var"
if I compile just the first TU without the foo () call in there, and
.global .align 4 .u32 var$lto_priv$1[1] = { 5 };
.global .align 4 .u32 var$lto_priv$0[1] = { 5 };
if I compile both, so there is some magic that makes these global and uglified
so that there can be multiple such vars.
No idea where that happens, whether it is NVPTX specific etc.
