https://gcc.gnu.org/bugzilla/show_bug.cgi?id=106447
Bug ID: 106447
Summary: nvptx offloading: C++ '__iterator_category' 'return'
ICE
Product: gcc
Version: 13.0
Status: UNCONFIRMED
Keywords: ice-on-valid-code, openmp
Severity: normal
Priority: P3
Component: target
Assignee: unassigned at gcc dot gnu.org
Reporter: tschwinge at gcc dot gnu.org
CC: sayle at gcc dot gnu.org, vries at gcc dot gnu.org
Target Milestone: ---
Target: nvptx
I found this during my "GPU support for minimal C++ library" investigations,
PR101544. You're not currently able to reproduce this, as it depends on a
number of WIP patches, but I'm filing it for later reference. (Or, I suppose
could maybe come up with a stand-alone reproducer.) I've done a quick try
reproducing the issue without offloading, with a nvptx target toolchain, but
didn't succeed to trigger the ICE there.
Simple program (similar to PR106445, but using a constructor with 'count'):
#include <vector>
int main()
{
#pragma omp target
{
std::vector<int> v(100);
}
return 0;
}
This explodes in nvptx offloading compilation:
during RTL pass: expand
dump file: ./a.xnvptx-none.mkoffload.254r.expand
[...]/build-gcc/x86_64-pc-linux-gnu/libstdc++-v3/include/bits/stl_iterator_base_types.h:
In function ‘__iterator_category’:
[...]/build-gcc/x86_64-pc-linux-gnu/libstdc++-v3/include/bits/stl_iterator_base_types.h:239:65:
internal compiler error: in emit_move_insn, at expr.cc:4039
239 | { return typename iterator_traits<_Iter>::iterator_category();
}
| ^
That's, through some layers of C++ templates,
'libstdc++-v3/include/bits/stl_iterator_base_types.h':
235 template<typename _Iter>
236 inline _GLIBCXX_CONSTEXPR
237 typename iterator_traits<_Iter>::iterator_category
238 __iterator_category(const _Iter&)
239 { return typename iterator_traits<_Iter>::iterator_category(); }
'*.253t.optimized':
;; Function __iterator_category
(_ZSt19__iterator_categoryIPiENSt15iterator_traitsIT_E17iterator_categoryERKS2_,
funcdef_no=39, decl_uid=1570, cgraph_uid=40, symbol_order=223)
__attribute__((omp declare target))
struct iterator_category __iterator_category (int * const & D.1815)
{
struct iterator_category D.1816;
<bb 2> :
<bb 3> :
gimple_label <<L0>>
gimple_return <D.1816>
}
'*.254r.expand':
;; Function __iterator_category
(_ZSt19__iterator_categoryIPiENSt15iterator_traitsIT_E17iterator_categoryERKS2_,
funcdef_no=39, decl_uid=1570, cgraph_uid=40, symbol_order=223)
;; Generating RTL for gimple basic block 2
;; Generating RTL for gimple basic block 3
EMERGENCY DUMP:
__attribute__((omp declare target))
struct iterator_category __iterator_category (int * const & D.1815)
{
int * const & _2(D);
(note 5 1 2 4 [bb 4] NOTE_INSN_BASIC_BLOCK)
(insn 2 5 3 4 (set (reg:DI 24)
(unspec:DI [
(const_int 0 [0])
] UNSPEC_ARG_REG))
"[...]/build-gcc/x86_64-pc-linux-gnu/libstdc++-v3/include/bits/stl_iterator_base_types.h":238:5
-1
(nil))
(insn 3 2 4 4 (set (mem/f/c:DI (reg/f:DI 17 virtual-stack-vars) [23
D.1815+0 S8 A64])
(reg:DI 24))
"[...]/build-gcc/x86_64-pc-linux-gnu/libstdc++-v3/include/bits/stl_iterator_base_types.h":238:5
-1
(nil))
(note 4 3 6 4 NOTE_INSN_FUNCTION_BEG)
(note 6 4 7 2 [bb 2] NOTE_INSN_BASIC_BLOCK)
(code_label 7 6 8 68 (nil) [0 uses])
(note 8 7 0 [bb 3] NOTE_INSN_BASIC_BLOCK)
}
I'm not attaching a nvptx offloading 'lto1' reproducer, as these are somewhat
old GCC sources with local modifications, so I'm not sure the LTO stream format
would match, etc. If someone is interested in reproducing this, we shall
coordinate.
(gdb) bt
#0 fancy_abort (file=file@entry=0x1927f68 "[...]/source-gcc/gcc/expr.cc",
line=line@entry=4039, function=function@entry=0x19298d8
<emit_move_insn(rtx_def*, rtx_def*)::__FUNCTION__> "emit_move_insn") at
[...]/source-gcc/gcc/diagnostic.cc:2012
#1 0x0000000000847b3f in emit_move_insn (x=x@entry=0x7ffff7690f60,
y=y@entry=0x7ffff7690f30) at [...]/source-gcc/gcc/expr.cc:4038
#2 0x00000000006e8119 in expand_value_return
(val=val@entry=0x7ffff7690f30) at [...]/source-gcc/gcc/cfgexpand.cc:3741
#3 0x00000000006e85b1 in expand_return
(retval=retval@entry=0x7ffff768abe0) at [...]/source-gcc/gcc/cfgexpand.cc:3813
#4 0x00000000006e8a1e in expand_gimple_stmt_1
(stmt=stmt@entry=0x7ffff76920f0) at [...]/source-gcc/gcc/cfgexpand.cc:3920
#5 0x00000000006e90da in expand_gimple_stmt
(stmt=stmt@entry=0x7ffff76920f0) at [...]/source-gcc/gcc/cfgexpand.cc:4046
#6 0x00000000006f3c88 in expand_gimple_basic_block
(bb=bb@entry=0x7ffff7687d68, disable_tail_calls=disable_tail_calls@entry=false)
at [...]/source-gcc/gcc/cfgexpand.cc:6093
#7 0x00000000006f6378 in (anonymous namespace)::pass_expand::execute
(this=0x1f98db0, fun=0x7ffff76250b8) at [...]/source-gcc/gcc/cfgexpand.cc:6819
[...]
(gdb) up
#1 0x0000000000847b3f in emit_move_insn (x=x@entry=0x7ffff7690f60,
y=y@entry=0x7ffff7690f30) at [...]/source-gcc/gcc/expr.cc:4038
4038 gcc_assert (mode != BLKmode
(gdb) list
4033 machine_mode mode = GET_MODE (x);
4034 rtx y_cst = NULL_RTX;
4035 rtx_insn *last_insn;
4036 rtx set;
4037
4038 gcc_assert (mode != BLKmode
4039 && (GET_MODE (y) == mode || GET_MODE (y) ==
VOIDmode));
4040
4041 /* If we have a copy that looks like one of the following
patterns:
4042 (set (subreg:M1 (reg:M2 ...)) (subreg:M1 (reg:M2 ...)))
(gdb) print mode
$1 = E_SImode
(gdb) print y
$2 = (rtx) 0x7ffff7690f30
(gdb) call debug_rtx(y)
(reg:QI 22 [ D.1816 ])
(gdb) call debug_rtx(x)
(reg:SI 23 [ <retval> ])
