https://gcc.gnu.org/bugzilla/show_bug.cgi?id=119485
Bug ID: 119485
Summary: OpenACC offloading compilation failure/ICE for C++
templated library functions
Product: gcc
Version: 15.0
Status: UNCONFIRMED
Keywords: openacc
Severity: normal
Priority: P3
Component: middle-end
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,
rguenth at gcc dot gnu.org
Blocks: 118518
Target Milestone: ---
Target: GCN, nvptx
For example, the following simple code, reduced from PR118518:
#include <cmath>
int main()
{
int r;
#pragma acc serial copyout(r)
{
r = (int) std::pow(1982, 0.495);
}
if (r != 42)
__builtin_abort();
}
For '-O0', GCN offloading:
ld: error: undefined symbol:
_ZSt3powIidEN9__gnu_cxx11__promote_2IDTplcvNS1_IT_XsrSt12__is_integerIS2_E7__valueEE6__typeELi0EcvNS1_IT0_XsrS3_IS7_E7__valueEE6__typeELi0EEXsrS3_ISB_E7__valueEE6__typeES2_S7_
>>> referenced by
./a.xamdgcn-amdhsa.mkoffload.hsaco-a.xamdgcn-amdhsa.mkoffload.2.o:(main._omp_fn.0)
>>> referenced by
./a.xamdgcn-amdhsa.mkoffload.hsaco-a.xamdgcn-amdhsa.mkoffload.2.o:(main._omp_fn.0)
collect2: error: ld returned 1 exit status
gcn mkoffload: fatal error:
build-gcc/gcc/x86_64-pc-linux-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit
status
For '-O0', nvptx offloading:
ptxas ./a.xnvptx-none.mkoffload.o, line 43; error : Call to
'_ZSt3powIidEN9__gnu_cxx11__promote_2IDTplcvNS1_IT_XsrSt12__is_integerIS2_E7__valueEE6__typeELi0EcvNS1_IT0_XsrS3_IS7_E7__valueEE6__typeELi0EEXsrS3_ISB_E7__valueEE6__typeES2_S7_'
requires call prototype
ptxas ./a.xnvptx-none.mkoffload.o, line 43; error : Unknown symbol
'_ZSt3powIidEN9__gnu_cxx11__promote_2IDTplcvNS1_IT_XsrSt12__is_integerIS2_E7__valueEE6__typeELi0EcvNS1_IT0_XsrS3_IS7_E7__valueEE6__typeELi0EEXsrS3_ISB_E7__valueEE6__typeES2_S7_'
ptxas fatal : Ptx assembly aborted due to errors
nvptx-as: ptxas returned 255 exit status
nvptx mkoffload: fatal error:
build-gcc/gcc/x86_64-pc-linux-gnu-accel-nvptx-none-gcc returned 1 exit status
That's (for both):
$ c++filt
_ZSt3powIidEN9__gnu_cxx11__promote_2IDTplcvNS1_IT_XsrSt12__is_integerIS2_E7__valueEE6__typeELi0EcvNS1_IT0_XsrS3_IS7_E7__valueEE6__typeELi0EEXsrS3_ISB_E7__valueEE6__typeES2_S7_
__gnu_cxx::__promote_2<decltype (((__gnu_cxx::__promote_2<int,
std::__is_integer<int>::__value>::__type)(0))+((__gnu_cxx::__promote_2<double,
std::__is_integer<double>::__value>::__type)(0))), std::__is_integer<decltype
(((__gnu_cxx::__promote_2<int,
std::__is_integer<int>::__value>::__type)(0))+((__gnu_cxx::__promote_2<double,
std::__is_integer<double>::__value>::__type)(0)))>::__value>::__type
std::pow<int, double>(int, double)
..., so some C++y masquerade of 'pow' (math power function).
Why in the nvptx case this doesn't result in the usual 'ld' 'unresolved symbol
[...]' error, might or might not be the same issue already reported/discussed
in PR106445 "nvptx offloading: C++ constructor symbol alias getting lost", PR
"[nvptx] Incorrect ptx code-gen for C++ code with templates", which I'm looking
into.
With optimizations enabled (and implicitly, inlining enabled), both GCN and
nvptx offloading compilation succeeds.
However, for '-O1 -fno-inline', for both GCN or nvptx offloading:
during IPA pass: inline
lto1: internal compiler error: Segmentation fault
Program received signal SIGSEGV, Segmentation fault.
0x0000000000c21c57 in offline_size (node=node@entry=0x7ffff76fa770,
info=info@entry=0x0) at ../../source-gcc/gcc/ipa-inline-analysis.cc:452
452 return info->size;
(gdb) bt
#0 0x0000000000c21c57 in offline_size (node=node@entry=0x7ffff76fa770,
info=info@entry=0x0) at ../../source-gcc/gcc/ipa-inline-analysis.cc:452
#1 0x0000000000c22064 in growth_positive_p
(node=node@entry=0x7ffff76fa770, known_edge=known_edge@entry=0x0,
edge_growth=edge_growth@entry=-2147483648) at
../../source-gcc/gcc/ipa-inline-analysis.cc:548
#2 0x000000000202850f in want_inline_function_to_all_callers_p
(cold=false, node=0x7ffff76fa770) at ../../source-gcc/gcc/ipa-inline.cc:1245
#3 ipa_inline () at ../../source-gcc/gcc/ipa-inline.cc:2956
#4 (anonymous namespace)::pass_ipa_inline::execute (this=<optimized out>)
at ../../source-gcc/gcc/ipa-inline.cc:3291
#5 0x0000000000dfa67b in execute_one_pass (pass=pass@entry=0x350b580) at
../../source-gcc/gcc/passes.cc:2659
#6 0x0000000000dfc6b2 in execute_ipa_pass_list (pass=0x350b580) at
../../source-gcc/gcc/passes.cc:3112
#7 0x000000000095ca3d in ipa_passes () at
../../source-gcc/gcc/cgraphunit.cc:2286
#8 symbol_table::compile (this=0x7ffff76ef000) at
../../source-gcc/gcc/cgraphunit.cc:2351
#9 0x000000000095db36 in symbol_table::compile (this=<optimized out>) at
../../source-gcc/gcc/cgraphunit.cc:2456
#10 0x0000000000871a2b in lto_main () at
../../source-gcc/gcc/lto/lto.cc:691
#11 0x0000000000f4903a in compile_file () at
../../source-gcc/gcc/toplev.cc:452
#12 0x000000000083b885 in do_compile () at
../../source-gcc/gcc/toplev.cc:2208
#13 toplev::main (this=this@entry=0x7fffffffd3ce, argc=argc@entry=17,
argv=0x349df60, argv@entry=0x7fffffffd4f8) at
../../source-gcc/gcc/toplev.cc:2371
#14 0x000000000083dc7b in main (argc=17, argv=0x7fffffffd4f8) at
../../source-gcc/gcc/main.cc:39
(gdb) call node->debug()
_ZSt3powIidEN9__gnu_cxx11__promote_2IDTplcvNS1_IT_XsrSt12__is_integerIS2_E7__valueEE6__typeELi0EcvNS1_IT0_XsrS3_IS7_E7__valueEE6__typeELi0EEXsrS3_ISB_E7__valueEE6__typeES2_S7_/2
(pow)
Type: function definition
Visibility: semantic_interposition prevailing_def_ironly
References:
Referring:
Read from file: a-pr118518-1-main_acc_.o
Availability: not_available
Unit id: 1
Function flags: count:1073741824 (estimated locally)
Called by: main._omp_fn.0/1 (1073741824 (estimated locally),1.00 per
call)
Calls:
That means, in 'gcc/ipa-inline-analysis.cc:growth_positive_p',
'ipa_size_summaries->get (node)' returned 'NULL', which means the 'node'
unexpectedly doesn't "exist"?
Conceptually similar (as far as I can tell; ICEs just in a different place) for
'-O2 -fno-inline' and higher, again for both GCN or nvptx offloading.
If I call specific 'std::powf', for example, instead of generic 'std::pow',
things work as expected for '-fno-inline'.
Standard library functions are intended to implicitly be tagged '#pragma acc
routine seq', and therefore be available for offloading use. So I suspect we
fail to do this here, in some way?
Referenced Bugs:
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=118518
[Bug 118518] gcc 14.2.1 nvptx cross compiler complains about alias definitions
in a struct with two constructors that are not aliases
