[Bug target/106447] New: nvptx offloading: C++ '__iterator_category' 'return' ICE

[Bug target/106447] New: nvptx offloading: C++ '__iterator_category' 'return' ICE

[tschwinge at gcc dot gnu.org]

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> ])

[Bug target/106447] nvptx offloading: C++ '__iterator_category' 'return' ICE

[roger at nextmovesoftware dot com]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=106447

Roger Sayle <roger at nextmovesoftware dot com> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
                 CC|                            |roger at nextmovesoftware dot com

--- Comment #1 from Roger Sayle <roger at nextmovesoftware dot com> ---
It looks like the problem is caused by the way QImode isn't a first class
citizen on nvptx, so QImode values are typically promoted internally to SImode.
 The problem will be in cfgexpand.cc's expand_value_return, which is the next
function up in the call stack, which has called emit_move_insn with mismatched
RTX modes.
It would be good to know what mode and old_mode are in this function, and
therefore why the call to convert_modes isn't doing what its supposed to.