https://gcc.gnu.org/bugzilla/show_bug.cgi?id=119692
Bug ID: 119692
Summary: C++ 'typeinfo', 'vtable' vs. OpenACC, OpenMP 'target'
offloading
Product: gcc
Version: 15.0
Status: UNCONFIRMED
Keywords: openacc, openmp
Severity: normal
Priority: P3
Component: c++
Assignee: unassigned at gcc dot gnu.org
Reporter: tschwinge at gcc dot gnu.org
CC: burnus at gcc dot gnu.org, jakub at gcc dot gnu.org
Target Milestone: ---
Created attachment 61052
--> https://gcc.gnu.org/bugzilla/attachment.cgi?id=61052&action=edit
'pr.C'
The attached 'pr.C' is meant to execute and terminate normally, but with
'-fopenmp' and offload device available, we get:
Thread 1 "a.out" received signal SIGSEGV, Segmentation fault.
0x00007ffff5e87655 in ?? () from /lib/x86_64-linux-gnu/libcuda.so.1
(gdb) bt
[...]
#8 0x00007ffff602a61e in ?? () from /lib/x86_64-linux-gnu/libcuda.so.1
#9 0x00007ffff79a398c in GOMP_OFFLOAD_dev2host (ord=<optimized out>,
dst=0x402020 <typeinfo for C2>, src=0x7fffcb600018, n=24) at
../../../source-gcc/libgomp/plugin/plugin-nvptx.c:2073
#10 0x00007ffff7c2db38 in gomp_device_copy (devicep=0x4faa40,
copy_func=<optimized out>, dst=0x7ffff7c442c9 "host", dstaddr=0x402020
<typeinfo for C2>, src=0x7ffff7c442e0 "dev", srcaddr=0x7fffcb600018, size=24)
at ../../../source-gcc/libgomp/target.c:247
#11 gomp_copy_dev2host (devicep=0x4faa40, aq=0x0, h=0x402020 <typeinfo for
C2>, d=0x7fffcb600018, sz=24) at ../../../source-gcc/libgomp/target.c:461
#12 gomp_unmap_vars_internal (tgt=0xb6ace0, do_copyfrom=true,
refcount_set=0x7fffffffd738, aq=0x0) at
../../../source-gcc/libgomp/target.c:2089
#13 gomp_unmap_vars (tgt=tgt@entry=0xb6ace0,
refcount_set=refcount_set@entry=0x7fffffffd738, do_copyfrom=true) at
../../../source-gcc/libgomp/target.c:2129
#14 0x00007ffff7c2f80f in GOMP_target_ext (device=<optimized out>,
fn=<optimized out>, mapnum=<optimized out>, hostaddrs=<optimized out>,
sizes=<optimized out>, kinds=<optimized out>, flags=<optimized out>,
depend=<optimized out>, args=<optimized out>) at
../../../source-gcc/libgomp/target.c:3307
#15 0x00000000004011e2 in main () at [...]/pr.C:13
Per '-fdump-tree-gimple', we've got :
#pragma omp target [...] map(tofrom:_ZTI2C2 [len: 24] [runtime_implicit])
map(tofrom:_ZTI2C1 [len: 16] [runtime_implicit]) map(tofrom:_ZTV2C1 [len: 24]
[runtime_implicit])
..., which are:
$ c++filt _ZTI2C2 _ZTI2C1 _ZTV2C1
typeinfo for C2
typeinfo for C1
vtable for C1
..., and per '*.s', they're in '.rodata'. (..., and the offloading machinery
wouldn't have any business modifying these, anyway?)
Why do we 'map' these 'tofrom' -- and how to fix this?
I suppose we should 'firstprivate' these, or at least only 'GOMP_MAP_TO'?
Would this be done in the gimplifier (via 'GOVD_FIRSTPRIVATE', or
'GOVD_MAP_TO_ONLY'?), or in the C++ front end (via langhook?)?
(Have not yet looked, but supposedly Fortran would do some similar handling for
its array descriptors?)
With OpenMP 'defaultmap(none)':
[...]/pr.C: In function ‘int main()’:
[...]/pr.C:16:8: error: ‘_ZTV2C1’ not specified in enclosing ‘target’
[...]/pr.C:13:9: note: enclosing ‘target’
[...]/pr.C:22:18: error: ‘_ZTI2C2’ not specified in enclosing ‘target’
[...]/pr.C:13:9: note: enclosing ‘target’
[...]/pr.C:22:18: error: ‘_ZTI2C1’ not specified in enclosing ‘target’
[...]/pr.C:13:9: note: enclosing ‘target’
Should 'defaultmap(none)' really report compiler-internal variables?
With OpenMP 'defaultmap(firstprivate)':
#pragma omp target [...] firstprivate(_ZTI2C2) firstprivate(_ZTI2C1)
firstprivate(_ZTV2C1)
..., and we get successful execution with offload device.
Similar for '-fopenacc':
#pragma omp target oacc_serial map(tofrom:_ZTI2C2 [len: 24])
map(tofrom:_ZTI2C1 [len: 16]) map(tofrom:_ZTV2C1 [len: 24])
..., but OpenACC 'default(none)' doesn't report these -- should it, like OpenMP
'defaultmap(none)'? That's probably because of 'DECL_ARTIFICIAL' handling in
'oacc_default_clause'.
(There's no OpenACC 'default(firstprivate)'.)
