[Bug other/90566] New: Support demangling with underscore-prefixed string after mangled name

[eyalroz at technion dot ac.il]

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

            Bug ID: 90566
           Summary: Support demangling with underscore-prefixed string
                    after mangled name
           Product: gcc
           Version: 6.3.0
            Status: UNCONFIRMED
          Severity: normal
          Priority: P3
         Component: other
          Assignee: unassigned at gcc dot gnu.org
          Reporter: eyalroz at technion dot ac.il
  Target Milestone: ---

libiberty performs the demangling for c++filt, the most commonly-used (and
perhaps only?) tool for demangling C++ names in object files and related file
formats. One of the "ecosystems" which produces such files is CUDA;
specifically in its intermediary representation for GPU code. 

Now, a GPU-device-side function, compiled with clang to PTX, can look like
this, for example: 

  .visible .entry _Z6squarePii(
          .param .u64 _Z6squarePii_param_0,
          .param .u32 _Z6squarePii_param_1
  )
  {

          ld.param.u32    %r1, [_Z6squarePii_param_1];
          mov.u32         %r2, %ctaid.x;
          setp.ge.s32     %p1, %r2, %r1;
          @%p1 bra        LBB6_2;
          ld.param.u64    %rd2, [_Z6squarePii_param_0];
          cvta.to.global.u64      %rd3, %rd2;
          mul.wide.u32    %rd4, %r2, 4;
          add.s64         %rd1, %rd3, %rd4;
          ld.global.u32   %r3, [%rd1];
          mul.lo.s32      %r4, %r3, %r3;
          st.global.u32   [%rd1], %r4;
          ret; 
  }

(see https://godbolt.org/z/GcDTVh for cland and nvcc output)

which clearly has mangled names. However, it seems the function parameter name
is somewhat malformed, or non-standard - being a mangled name, followed
immediately by an underscore and more text: mangledblahblah_param_0.

When demangling, the function name gets demangled fine, but the parameter does
not:

    .visible .entry square(int*, int)(
            .param .u64 _Z6squarePii_param_0,
            .param .u32 _Z6squarePii_param_1
    )

and from what the c++filt people say - this is libiberty's output. I ask that
libiberty either auto-detect this case, or have an option to detect it; and
when that's turned on, demangle the above into:

  .visible .entry square(int*, int)(
          .param .u64 square(int*, int)_param_0,
          .param .u32 square(int*, int)_param_1
  )

or

  .visible .entry square(int*, int)(
          .param .u64 square(int*, int) param_0,
          .param .u32 square(int*, int) param_1
  )

or something else that's meaningful.

Caveat: I realize that libiberty is FOSS and CUDA involves a bunch of
closed-source software by a company notorious for keeping code and specs
closed, and not making it easy for FOSS developers. Still, we are talking about
something clang compiles; and it's only being mindful of an underscore.

(Note: I first filed this as a bug against c++filt:
https://sourceware.org/bugzilla/show_bug.cgi?id=24557 ) and was directed to
file here.