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

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

[Thomas Schwinge]

[-- 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 --]

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

[Nathan Sidwell]

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

[nvptx] fix some c++ tests

[Nathan Sidwell]

[-- 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

Re: [nvptx] fix some c++ tests

[Thomas Schwinge]

[-- 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 --]

Re: [nvptx] fix some c++ tests

[Bernd Schmidt]

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