https://gcc.gnu.org/bugzilla/show_bug.cgi?id=121573
--- Comment #2 from Benjamin Schulz <schulz.benjamin at googlemail dot com> ---
ok, yes.. interestingly, if i write the pragma following a variable
declaration, nvc++ accepts it too.
but it then, well creates two variable versions.
#include <openacc.h>
#include <stdio.h>
typedef struct
{
int a;
float b;
} mystruct;
mystruct dev_var={42, 3.14f};
#pragma acc declare device_resident (dev_var)
int main(void)
{
printf("on host: a=%d b=%f\n", dev_var.a, dev_var.b);
// Try to read back into a host variable
mystruct out;
#pragma acc parallel copyout(out)
{
if (acc_on_device(acc_device_nvidia)!=0)
printf("on nvidia\n");
else
printf("not on nvidia\n");
printf("on parallel region, should be 42; dev_var.a=%d \n", dev_var.a);
printf("on parallel region, should be 3.14: dev_var.b=%f \n",
dev_var.b);
out = dev_var;
}
printf("Device-resident struct copied back, should be 42 and 3.14: a=%d
b=%f\n", out.a, out.b);
return 0;
}
Works entirely on gcc....
The spec says:
"The device_resident clause specifies that the memory for the named variables
should be allocated in the current device memory and not in local memory"
in this example, i can initialize and access it from host...
What is interesting: nvc++ does not initialize it on device then;
on host: a=42 b=3.140000
on nvidia
on parallel region, should be 42; dev_var.a=0
on parallel region, should be 3.14: dev_var.b=0.000000
Device-resident struct copied back, should be 42 and 3.14: a=0 b=0.000000
gcc output of that snippet is this:
on host: a=42 b=3.140000
on nvidia
on parallel region, should be 42; dev_var.a=42
on parallel region, should be 3.14: dev_var.b=3.140000
Device-resident struct copied back, should be 42 and 3.14: a=42 b=3.140000
Process returned 0 (0x0) execution time : 0.229 s
Press ENTER to continue.
So gcc creates a variable on both host and device and initializes it...
hm
