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
pgpSvhWMBpv_L.pgp
Description: PGP signature
