[Bug binutils/24557] New: Support underscore-prefixed string after mangled name (for CUDA PTX parameters)

[eyalroz at technion dot ac.il]

https://sourceware.org/bugzilla/show_bug.cgi?id=24557

            Bug ID: 24557
           Summary: Support underscore-prefixed string after mangled name
                    (for CUDA PTX parameters)
           Product: binutils
           Version: 2.28
            Status: UNCONFIRMED
          Severity: enhancement
          Priority: P2
         Component: binutils
          Assignee: unassigned at sourceware dot org
          Reporter: eyalroz at technion dot ac.il
  Target Milestone: ---

c++filt is the most common (? 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 to PTX, can look like this, for
example:

  .visible .entry _ZN7kernels15determine_sizesILj0EEEvP7sizes_t(
      .param .u64 _ZN7kernels15determine_sizesILj0EEEvP7sizes_t_param_0
  )
  {
      .reg .b32   %r<6>;
      .reg .b64   %rd<3>;


     ld.param.u64    %rd1, 
[_ZN7kernels15determine_sizesILj0EEEvP7sizes_t_param_0];
      cvta.to.global.u64  %rd2, %rd1;
      mov.u32%r1, %dynamic_smem_size;
      st.global.u32   [%rd2], %r1;
      mov.u32%r2, %total_smem_size;
      mov.u32%r3, %dynamic_smem_size;
      sub.s32     %r5, %r2, %r3;
      st.global.u32   [%rd2+4], %r5;
      mov.u32%r4, %total_smem_size;
      st.global.u32   [%rd2+8], %r4;
      ret;
  }

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 void kernels::determine_sizes<0u>(sizes_t*)(
          .param .u64 _ZN7kernels15determine_sizesILj0EEEvP7sizes_t_param_0
  )

I would like c++filt to either auto-detect this case, or have an option to
detect it; and when that's turned on, demangle the above into:

  .visible .entry void kernels::determine_sizes<0u>(sizes_t*)(
          .param .u64 kernels::determine_sizes<0u>(sizes_t*)_param_0
  )

or

  .visible .entry void kernels::determine_sizes<0u>(sizes_t*)(
          .param .u64 kernels::determine_sizes<0u>(sizes_t*) param_0
  )

or something else that's meaningful.


Caveat: I realize that c++filt is FOSS and CUDA is closed-source software by a
company notorious for keeping code and specs closed, and not making it easy for
FOSS developers. So I wonder understand a "principled" objection to this.
Still, it's just an underscore.

-- 
You are receiving this mail because:
You are on the CC list for the bug.