https://gcc.gnu.org/bugzilla/show_bug.cgi?id=119692
--- Comment #4 from Tobias Burnus <burnus at gcc dot gnu.org> ---
Some generic remarks - regarding OpenMP to:
> When still 'map'ping C++ 'typeinfo', 'vtable' at the OpenACC compute,
> OpenMP 'target' construct
Obviously, the type info etc. is required for all work on that class on
a device side, like:
MyClass *x = new MyClass();
MyClass *y = new InheritFromMyClass();
However, on the OpenMP side, there are the following
RESTRICTIONS related to MAPPING:
* The run-time type information (RTTI) of an object can only be accessed
from the device on which it was constructed.
* Invoking a virtual member function of an object on a device other
than the device on which the object was constructed results in unspecified
behavior, unless the object is accessible and was constructed on the host
device.
The latter relates to the following; assume that 'func' is a virtual function
of the base class:
#pragma omp requires self_maps /* or/and */ unified_shared_memory
...
MyClass *y = new InheritFromMyClass();
#pragma omp target
{
y->func (123);
}
where 'y->vtable' is the host vtable and, hence, 'vtable->func' points to the
host function.
Similar to 'omp declare target(some_func) indirect' and using device function
pointers, the compiler has to insert the lookup, i.e.
the user's
y->func (123);
becomes internally
(GOMP_target_map_indirect_ptr (y->vtable->func) (123);
The function returns the device ptr, if the mapping has been found,
and otherwise the original pointer. Thus, it only works if the
function is 'declare target'; for virtual functions, the 'indirect'
clause is then implied, for other functions the 'indirect' is must
be specified by the user.
Regarding mainline support for USM, see 'unified_shared_memory' entry at
https://gcc.gnu.org/onlinedocs/libgomp/Offload-Target-Specifics.html
And note that the automatic insertion of the GOMP_target_map_indirect_ptr call
for virtual functions is an OpenMP 5.2 feature, which is not yet implemented in
GCC 15.
And regarding the example above: If actual mapping happens (i.e. no self map)
and the vtable exists on the device, it is of course permitted to update the
vtable pointer in the mapped class instance to the device's one. But that's not
required.
* * *
OpenMP 6.1 will (tentatively) add support for mapping polymorphic types in
Fortran, starting with having the Fortran's vtable also accessible (= USM
case).
