[nvptx] Incorrect %retval usage in C++ code

[Thomas Schwinge <thomas at codesourcery dot com>]

Hi!

I observed that for a (slowly increasing?) number of C++ testcases
(gcc/testsuite/g++.*/), their nvptx compilation fails as follows:

    spawn [...]/build-gcc/gcc/testsuite/g++/../../xg++ -B[...]/build-gcc/gcc/testsuite/g++/../../ [...]/source-gcc/gcc/testsuite/g++.dg/cpp0x/nsdmi4.C -fno-diagnostics-show-caret -fdiagnostics-color=never --sysroot=[...]/install/nvptx-none -fmessage-length=0 -std=c++11 -pedantic-errors -Wno-long-long -DNO_LABEL_VALUES -DNO_TRAMPOLINES -isystem [...]/build-gcc/nvptx-none/./newlib/targ-include -isystem [...]/source-gcc/newlib/libc/include -B[...]/build-gcc/nvptx-none/./newlib/ -L[...]/build-gcc/nvptx-none/./newlib -mmainkernel -lm -o ./nsdmi4.exe
    ptxas /tmp/cclwz6Zm.o, line 52; error   : Arguments mismatch for instruction 'mov'
    ptxas /tmp/cclwz6Zm.o, line 52; error   : Unknown symbol '%retval'
    ptxas /tmp/cclwz6Zm.o, line 52; error   : Label expected for forward reference of '%retval'
    ptxas fatal   : Ptx assembly aborted due to errors
    nvptx-as: ptxas returned 255 exit status
    compiler exited with status 1

Reduced from g++.dg/cpp0x/nsdmi4.C:

    $ < source-gcc/gcc/testsuite/g++.dg/cpp0x/nsdmi4_.C cat
    struct A
    {
      A() { }
      A(const A&) { }
    };
    
    A f() { return A(); }
    $ build-gcc/gcc/xg++ -Bbuild-gcc/gcc/ source-gcc/gcc/testsuite/g++.dg/cpp0x/nsdmi4_.C -std=c++11 -S
    $ < nsdmi4_.s c++filt
    // BEGIN PREAMBLE
            .version        3.1
            .target sm_30
            .address_size 64
    // END PREAMBLE
    
    // BEGIN FUNCTION DECL: A::A()
    .func A::A()(.param.u64 %in_ar1);
    // BEGIN FUNCTION DEF: A::A()
    .func A::A()(.param.u64 %in_ar1)
    {
            .reg.u64 %ar1;
            .reg.u64 %hr10;
            .reg.u64 %r22;
            .reg.u64 %frame;
            .local.align 8 .b8 %farray[8];
            cvta.local.u64 %frame, %farray;
            ld.param.u64 %ar1, [%in_ar1];
                    mov.u64 %r22, %ar1;
                    st.u64  [%frame], %r22;
            ret;
            }
    // BEGIN GLOBAL FUNCTION DECL: f()
    .visible .func f()(.param.u64 %in_ar1);
    // BEGIN GLOBAL FUNCTION DEF: f()
    .visible .func f()(.param.u64 %in_ar1)
    {
            .reg.u64 %ar1;
            .reg.u64 %hr10;
            .reg.u64 %r22;
            .reg.u64 %r23;
    ld.param.u64 %ar1, [%in_ar1];
                    mov.u64 %r22, %ar1;
                    mov.u64 %r23, %r22;
            {
                    .param.u64 %out_arg0;
                    st.param.u64 [%out_arg0], %r23;
                    call A::A(), (%out_arg0);
            }
                    mov.u64 %retval, %r22;
            ret;
            }

Notice the stray %retval usage very near the end.


Note that before r226901, Â[PR64164] Drop copyrename, use coalescible
partition as base when optimizing.Â,
<http://news.gmane.org/find-root.php?message_id=%3Cor37zlpujd.fsf%40livre.home%3E>,
this test case did set up a %frame, and did not assign to %retval:

    @@ -31,21 +31,18 @@
            .reg.u32 %r23;
            .reg.u64 %r24;
            .reg.u64 %r25;
    -       .reg.u64 %frame;
    -       .local.align 8 .b8 %farray[8];
    -       cvta.local.u64 %frame, %farray;
     ld.param.u64 %ar1, [%in_ar1];
                    mov.u64 %r24, %ar1;
    -               st.u64  [%frame], %r24;
                    ld.global.u32   %r22, [c];
                    add.u32 %r23, %r22, 1;
                    st.global.u32   [c], %r23;
    -               ld.u64  %r25, [%frame];
    +               mov.u64 %r25, %r24;
            {
                    .param.u64 %out_arg0;
                    st.param.u64 [%out_arg0], %r25;
                    call _ZN1AC1Ev, (%out_arg0);
            }
    +               mov.u64 %retval, %r24;
            ret;
            }

But I don't think that this recent commit is directly related to the
problem at hand, but it just exposes it some more: before that recent
commit, there have already been test cases failing with the same stray
%retval usage.

I suspect that this mov is incorrectly generated by
gcc/function.c:expand_function_end, but can't tell if something's wrong
in there, or rather in the nvptx backend's NVPTX_RETURN_REGNUM handling
(which I can't claim to really understand), and as I'm unlikely to spend
more time on this before leaving for vacations soon, I wanted to dump my
state now.  Maybe one of you has an idea about this.


Also, I guess the following cleanup (untested) is in order:

diff --git gcc/config/nvptx/nvptx.h gcc/config/nvptx/nvptx.h
index afe4fcd..d846ec3 100644
--- gcc/config/nvptx/nvptx.h
+++ gcc/config/nvptx/nvptx.h
@@ -103,8 +103,8 @@ enum reg_class
 #define N_REG_CLASSES (int) LIM_REG_CLASSES
 
 #define REG_CLASS_NAMES {	  \
-    "RETURN_REG",		  \
     "NO_REGS",			  \
+    "RETURN_REG",		  \
     "ALL_REGS" }
 
 #define REG_CLASS_CONTENTS	\
@@ -119,7 +119,7 @@ enum reg_class
 
 #define GENERAL_REGS ALL_REGS
 
-#define REGNO_REG_CLASS(R) ((R) == 4 ? RETURN_REG : ALL_REGS)
+#define REGNO_REG_CLASS(R) ((R) == NVPTX_RETURN_REGNUM ? RETURN_REG : ALL_REGS)
 
 #define BASE_REG_CLASS ALL_REGS
 #define INDEX_REG_CLASS NO_REGS


GrÃÃe,
 Thomas

Attachment: pgp1upIDhsNPb.pgp
Description: PGP signature