* [nvptx] Incorrect %retval usage in C++ code
@ 2015-08-17 7:37 Thomas Schwinge
2015-08-18 14:07 ` Nathan Sidwell
0 siblings, 1 reply; 5+ messages in thread
From: Thomas Schwinge @ 2015-08-17 7:37 UTC (permalink / raw)
To: Bernd Schmidt, Nathan Sidwell; +Cc: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 5303 bytes --]
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 #2: Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 5+ messages in thread
* Re: [nvptx] Incorrect %retval usage in C++ code
2015-08-17 7:37 [nvptx] Incorrect %retval usage in C++ code Thomas Schwinge
@ 2015-08-18 14:07 ` Nathan Sidwell
2015-10-08 17:34 ` [nvptx] fix some c++ tests Nathan Sidwell
0 siblings, 1 reply; 5+ messages in thread
From: Nathan Sidwell @ 2015-08-18 14:07 UTC (permalink / raw)
To: Thomas Schwinge, Bernd Schmidt; +Cc: gcc-patches
On 08/17/15 03:31, Thomas Schwinge wrote:
> Hi!
>
> I observed that for a (slowly increasing?) number of C++ testcases
> (gcc/testsuite/g++.*/), their nvptx compilation fails as follows:
This is occurring because in C++land (some?) functions that return structs via a
hidden parameter now also return a reference to that struct (see cp_genericize
in cp-gimplify.c). Hence it is possible for RETURN_IN_REG (return_type) to be
false AND for a return value in a register (!).
testing a patch.
nathan
^ permalink raw reply [flat|nested] 5+ messages in thread
* [nvptx] fix some c++ tests
@ 2015-10-08 17:34 ` Nathan Sidwell
2015-10-09 5:55 ` Thomas Schwinge
0 siblings, 1 reply; 5+ messages in thread
From: Nathan Sidwell @ 2015-10-08 17:34 UTC (permalink / raw)
To: GCC Patches; +Cc: Thomas Schwinge
[-- Attachment #1: Type: text/plain, Size: 228 bytes --]
I've committed this to trunk. The C++ ABI now returns a pointer to the
passed-in artificial arg that points to the return area. consequently
return-in-mem and type_mode(return_type) == VOIDmode are not tautologies.
nathan
[-- Attachment #2: trunk-ptx-cxx.patch --]
[-- Type: text/x-patch, Size: 2833 bytes --]
2015-10-08 Nathan Sidwell <nathan@acm.org>
* config/nvptx/nvptx.h (struct machine_function): Add comment.
* config/nvptx/nvptx.c (nvptx_declare_function_name): Functions
may return pointer as well as in memory.
(nvptx_output_return): Likewise.
Index: gcc/config/nvptx/nvptx.c
===================================================================
--- gcc/config/nvptx/nvptx.c (revision 228617)
+++ gcc/config/nvptx/nvptx.c (working copy)
@@ -531,13 +531,8 @@ nvptx_declare_function_name (FILE *file,
nvptx_write_function_decl (s, name, decl);
fprintf (file, "%s", s.str().c_str());
- bool return_in_mem = false;
- if (TYPE_MODE (result_type) != VOIDmode)
- {
- machine_mode mode = TYPE_MODE (result_type);
- if (!RETURN_IN_REG_P (mode))
- return_in_mem = true;
- }
+ bool return_in_mem = (TYPE_MODE (result_type) != VOIDmode
+ && !RETURN_IN_REG_P (TYPE_MODE (result_type)));
fprintf (file, "\n{\n");
@@ -547,9 +542,13 @@ nvptx_declare_function_name (FILE *file,
false, return_in_mem);
if (return_in_mem)
fprintf (file, "\t.reg.u%d %%ar1;\n", GET_MODE_BITSIZE (Pmode));
- else if (TYPE_MODE (result_type) != VOIDmode)
+
+ /* C++11 ABI causes us to return a reference to the passed in
+ pointer for return_in_mem. */
+ if (cfun->machine->ret_reg_mode != VOIDmode)
{
- machine_mode mode = arg_promotion (TYPE_MODE (result_type));
+ machine_mode mode = arg_promotion
+ ((machine_mode)cfun->machine->ret_reg_mode);
fprintf (file, "\t.reg%s %%retval;\n",
nvptx_ptx_type_from_mode (mode, false));
}
@@ -635,17 +634,13 @@ nvptx_declare_function_name (FILE *file,
const char *
nvptx_output_return (void)
{
- tree fntype = TREE_TYPE (current_function_decl);
- tree result_type = TREE_TYPE (fntype);
- if (TYPE_MODE (result_type) != VOIDmode)
+ machine_mode mode = (machine_mode)cfun->machine->ret_reg_mode;
+
+ if (mode != VOIDmode)
{
- machine_mode mode = TYPE_MODE (result_type);
- if (RETURN_IN_REG_P (mode))
- {
- mode = arg_promotion (mode);
- fprintf (asm_out_file, "\tst.param%s\t[%%out_retval], %%retval;\n",
- nvptx_ptx_type_from_mode (mode, false));
- }
+ mode = arg_promotion (mode);
+ fprintf (asm_out_file, "\tst.param%s\t[%%out_retval], %%retval;\n",
+ nvptx_ptx_type_from_mode (mode, false));
}
return "ret;";
Index: gcc/config/nvptx/nvptx.h
===================================================================
--- gcc/config/nvptx/nvptx.h (revision 228617)
+++ gcc/config/nvptx/nvptx.h (working copy)
@@ -228,7 +228,7 @@ struct GTY(()) machine_function
bool has_call_with_varargs;
bool has_call_with_sc;
HOST_WIDE_INT outgoing_stdarg_size;
- int ret_reg_mode;
+ int ret_reg_mode; /* machine_mode not defined yet. */
int punning_buffer_size;
};
#endif
^ permalink raw reply [flat|nested] 5+ messages in thread
* Re: [nvptx] fix some c++ tests
2015-10-08 17:34 ` [nvptx] fix some c++ tests Nathan Sidwell
@ 2015-10-09 5:55 ` Thomas Schwinge
2015-10-09 10:19 ` Bernd Schmidt
0 siblings, 1 reply; 5+ messages in thread
From: Thomas Schwinge @ 2015-10-09 5:55 UTC (permalink / raw)
To: Nathan Sidwell; +Cc: GCC Patches
[-- Attachment #1: Type: text/plain, Size: 5336 bytes --]
Hi Nathan!
Thanks for looking into this. When I reported this,
<http://news.gmane.org/find-root.php?message_id=%3C87fv3is74z.fsf%40schwinge.name%3E>,
quite a lot of testcases had been failing -- with recent GCC trunk, the
number is smaller (because of other middle end/optimization changes, I
suppose). Yet, the problem still can be observed; grep for "%retval" in
g++.log.
On Thu, 8 Oct 2015 13:33:55 -0400, Nathan Sidwell <nathan@acm.org> wrote:
> I've committed this to trunk. The C++ ABI now returns a pointer to the
> passed-in artificial arg that points to the return area. consequently
> return-in-mem and type_mode(return_type) == VOIDmode are not tautologies.
> * config/nvptx/nvptx.c (nvptx_declare_function_name): Functions
> may return pointer as well as in memory.
> (nvptx_output_return): Likewise.
Hmm, but what I now see happening is that now there is incorrect
"%out_retval" usage (but it's not part of the function prototypes), for
example:
$ build-gcc/gcc/xg++ -Bbuild-gcc/gcc/ source-gcc/gcc/testsuite/g++.dg/ipa/ipa-icf-6.C -std=gnu++98 -O3 -o ipa-icf-6.exe -S
Diff before/after your change:
$ diff -U4 O/ipa-icf-6.exe ipa-icf-6.exe
--- O/ipa-icf-6.exe 2015-10-06 18:30:21.581187448 +0200
+++ ipa-icf-6.exe 2015-10-09 07:49:23.894893180 +0200
@@ -10,8 +10,9 @@
.visible .func _Z5test21A(.param.u64 %in_ar1, .param.u64 %in_ar2)
{
.reg.u64 %ar2;
.reg.u64 %ar1;
+ .reg.u64 %retval;
.reg.u64 %hr10;
.reg.u64 %r25;
.reg.u64 %r26;
.reg.u64 %r27;
@@ -28,8 +29,9 @@
}
$L2:
st.u64 [%r25+8], %r25;
mov.u64 %retval, %r25;
+ st.param.u64 [%out_retval], %retval;
ret;
}
// BEGIN FUNCTION DECL: _ZL7test_mePF1AS_E
.func _ZL7test_mePF1AS_E(.param.u64 %in_ar1);
@@ -78,8 +80,9 @@
.visible .func _Z5test11A(.param.u64 %in_ar1, .param.u64 %in_ar2)
{
.reg.u64 %ar2;
.reg.u64 %ar1;
+ .reg.u64 %retval;
.reg.u64 %hr10;
.reg.u64 %r25;
.reg.u64 %r26;
.reg.u64 %r27;
@@ -96,8 +99,9 @@
}
$L6:
st.u64 [%r25+8], %r25;
mov.u64 %retval, %r25;
+ st.param.u64 [%out_retval], %retval;
ret;
}
// BEGIN GLOBAL FUNCTION DECL: main
.visible .func (.param.u32 %out_retval)main(.param.u32 %argc, .param.u64 %argv);
(I have not yet made an attempt at trying to understand the problem.)
For reference:
> Index: gcc/config/nvptx/nvptx.c
> ===================================================================
> --- gcc/config/nvptx/nvptx.c (revision 228617)
> +++ gcc/config/nvptx/nvptx.c (working copy)
> @@ -531,13 +531,8 @@ nvptx_declare_function_name (FILE *file,
> nvptx_write_function_decl (s, name, decl);
> fprintf (file, "%s", s.str().c_str());
>
> - bool return_in_mem = false;
> - if (TYPE_MODE (result_type) != VOIDmode)
> - {
> - machine_mode mode = TYPE_MODE (result_type);
> - if (!RETURN_IN_REG_P (mode))
> - return_in_mem = true;
> - }
> + bool return_in_mem = (TYPE_MODE (result_type) != VOIDmode
> + && !RETURN_IN_REG_P (TYPE_MODE (result_type)));
>
> fprintf (file, "\n{\n");
>
> @@ -547,9 +542,13 @@ nvptx_declare_function_name (FILE *file,
> false, return_in_mem);
> if (return_in_mem)
> fprintf (file, "\t.reg.u%d %%ar1;\n", GET_MODE_BITSIZE (Pmode));
> - else if (TYPE_MODE (result_type) != VOIDmode)
> +
> + /* C++11 ABI causes us to return a reference to the passed in
> + pointer for return_in_mem. */
> + if (cfun->machine->ret_reg_mode != VOIDmode)
> {
> - machine_mode mode = arg_promotion (TYPE_MODE (result_type));
> + machine_mode mode = arg_promotion
> + ((machine_mode)cfun->machine->ret_reg_mode);
> fprintf (file, "\t.reg%s %%retval;\n",
> nvptx_ptx_type_from_mode (mode, false));
> }
> @@ -635,17 +634,13 @@ nvptx_declare_function_name (FILE *file,
> const char *
> nvptx_output_return (void)
> {
> - tree fntype = TREE_TYPE (current_function_decl);
> - tree result_type = TREE_TYPE (fntype);
> - if (TYPE_MODE (result_type) != VOIDmode)
> + machine_mode mode = (machine_mode)cfun->machine->ret_reg_mode;
> +
> + if (mode != VOIDmode)
> {
> - machine_mode mode = TYPE_MODE (result_type);
> - if (RETURN_IN_REG_P (mode))
> - {
> - mode = arg_promotion (mode);
> - fprintf (asm_out_file, "\tst.param%s\t[%%out_retval], %%retval;\n",
> - nvptx_ptx_type_from_mode (mode, false));
> - }
> + mode = arg_promotion (mode);
> + fprintf (asm_out_file, "\tst.param%s\t[%%out_retval], %%retval;\n",
> + nvptx_ptx_type_from_mode (mode, false));
> }
>
> return "ret;";
> Index: gcc/config/nvptx/nvptx.h
> ===================================================================
> --- gcc/config/nvptx/nvptx.h (revision 228617)
> +++ gcc/config/nvptx/nvptx.h (working copy)
> @@ -228,7 +228,7 @@ struct GTY(()) machine_function
> bool has_call_with_varargs;
> bool has_call_with_sc;
> HOST_WIDE_INT outgoing_stdarg_size;
> - int ret_reg_mode;
> + int ret_reg_mode; /* machine_mode not defined yet. */
> int punning_buffer_size;
> };
> #endif
Grüße,
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 5+ messages in thread
* Re: [nvptx] fix some c++ tests
2015-10-09 5:55 ` Thomas Schwinge
@ 2015-10-09 10:19 ` Bernd Schmidt
0 siblings, 0 replies; 5+ messages in thread
From: Bernd Schmidt @ 2015-10-09 10:19 UTC (permalink / raw)
To: Thomas Schwinge, Nathan Sidwell; +Cc: GCC Patches
On 10/09/2015 07:55 AM, Thomas Schwinge wrote:
> Hi Nathan!
>
> Thanks for looking into this. When I reported this,
> <http://news.gmane.org/find-root.php?message_id=%3C87fv3is74z.fsf%40schwinge.name%3E>,
> quite a lot of testcases had been failing -- with recent GCC trunk, the
> number is smaller (because of other middle end/optimization changes, I
> suppose). Yet, the problem still can be observed; grep for "%retval" in
> g++.log.
>
> On Thu, 8 Oct 2015 13:33:55 -0400, Nathan Sidwell <nathan@acm.org> wrote:
>> I've committed this to trunk. The C++ ABI now returns a pointer to the
>> passed-in artificial arg that points to the return area. consequently
>> return-in-mem and type_mode(return_type) == VOIDmode are not tautologies.
>
>> * config/nvptx/nvptx.c (nvptx_declare_function_name): Functions
>> may return pointer as well as in memory.
>> (nvptx_output_return): Likewise.
>
> Hmm, but what I now see happening is that now there is incorrect
> "%out_retval" usage (but it's not part of the function prototypes), for
> example:
Does the caller end up using that extra pointer? Otherwise, just reusing
TARGET_OMIT_STRUCT_RETURN_REG might be useful.
Bernd
^ permalink raw reply [flat|nested] 5+ messages in thread
end of thread, other threads:[~2015-10-09 10:19 UTC | newest]
Thread overview: 5+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2015-08-17 7:37 [nvptx] Incorrect %retval usage in C++ code Thomas Schwinge
2015-08-18 14:07 ` Nathan Sidwell
2015-10-08 17:34 ` [nvptx] fix some c++ tests Nathan Sidwell
2015-10-09 5:55 ` Thomas Schwinge
2015-10-09 10:19 ` Bernd Schmidt
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).