https://gcc.gnu.org/bugzilla/show_bug.cgi?id=121170
Bug ID: 121170
Summary: gcc refuses to compile if a struct member variable
appears in an OpenAcc construct, which nvc++ accepts
Product: gcc
Version: 15.1.0
Status: UNCONFIRMED
Severity: normal
Priority: P3
Component: target
Assignee: unassigned at gcc dot gnu.org
Reporter: schulz.benjamin at googlemail dot com
Target Milestone: ---
In contrast to OpenMP, which seems to forbid to use member variables of structs
and classes in constructs unless one calls the openmp loop in these, see
https://www.openmp.org/spec-html/5.2/openmpsu19.html
Nvidia's OpenAcc does not appear to have such (strange and often performance
limiting) restrictions.
According to
https://www.openacc.org/sites/default/files/inline-images/Specification/OpenACC-3.3-final.pdf
isdeviceptr only needs a pointer. There is nothing in the standard that forbids
that this pointer is a struct or class member
Compile the following test case:
#include <openacc.h>
struct mytensor
{
int *strides;
int *extents;
double *data;
};
int main()
{
mytensor t;
t.data=(double*)acc_malloc(sizeof(double)*20);
#pragma acc parallel loop gang deviceptr(t.data)
for(int i=1;i<20;i++)
{
t.data[i]=20;
}
}
with nvc++ ./main.cpp -acc -gpu=cuda12.9
nvc++ will compile it without warning and generate valid cuda code.
In contrast, gcc will stop compilation, saying:
/home/benni/projects/openmptestnew/openmpoffloatest/main.cpp:16:50: Error:
expected »)« before ».« token
16 | #pragma acc parallel loop gang deviceptr(t.data)
| ^
| )
/home/benni/projects/openmptestnew/openmpoffloatest/main.cpp:16:49: Error: »t«
is not a pointer variable
16 | #pragma acc parallel loop gang deviceptr(t.data)
| ^
well double *x=t.data needs no conversion, so t.data is a pointer variable.
Also Nvidia's compiler generates valid cuda code from it, and as Nvidia created
the OpenAcc standard, I would suggest what their compiler accepts is the
working reference for that.
(The fact that OpenMP refuses to accept struct members in constructs may create
performance overhead. Since, for example, one can then not declare such struct
members as shared in OpenMP creating unnecessary copies, or preventing to use
structs to order the code)
