[Bug target/104815] New: [nvptx] Use bitbucket operand when REG_UNUSED

[Bug target/104815] New: [nvptx] Use bitbucket operand when REG_UNUSED

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104815

            Bug ID: 104815
           Summary: [nvptx] Use bitbucket operand when REG_UNUSED
           Product: gcc
           Version: 12.0
            Status: UNCONFIRMED
          Severity: enhancement
          Priority: P3
         Component: target
          Assignee: unassigned at gcc dot gnu.org
          Reporter: vries at gcc dot gnu.org
  Target Milestone: ---

Consider this source code:
...
enum memmodel
{
  MEMMODEL_RELAXED = 0
};

unsigned long long int *p64;
unsigned long long int v64;

int
main()
{
  __atomic_fetch_add (p64, v64, MEMMODEL_RELAXED);

  return 0;
}
...

It results in this code:
...
{
        .reg.u32 %value;
        .reg.u64 %r25;
        .reg.u64 %r26;
        .reg.u64 %r27;
                ld.global.u64   %r25, [p64];
                ld.global.u64   %r27, [v64];
                atom.add.u64    %r26, [%r25], %r27;
                mov.u32 %value, 0;
        st.param.u32    [%value_out], %value;
        ret;
}
...

We can however do:
...
                atom.add.u64    _, [%r25], %r27;
...
because %r26 is not used anywhere.

We can detect this situation in the insn by finding a reg-note:
...
/* Identifies a register set in this insn and never used.  */
REG_NOTE (UNUSED)
...

Currently testing this code:
...
diff --git a/gcc/config/nvptx/nvptx.cc b/gcc/config/nvptx/nvptx.cc
index c6cec0c27c22..c7223737cca5 100644
--- a/gcc/config/nvptx/nvptx.cc
+++ b/gcc/config/nvptx/nvptx.cc
@@ -538,7 +538,15 @@ output_reg (FILE *file, unsigned regno, machine_mode
inner_mode,

       if (HARD_REGISTER_NUM_P (regno))
        fprintf (file, "%s", reg_names[regno]);
       else
-       fprintf (file, "%%r%d", regno);
+       {
+         rtx reg = regno_reg_rtx[regno];
+         if (current_output_insn
+             && (find_reg_note (current_output_insn, REG_UNUSED, reg)
+                 != NULL_RTX))
+           fprintf (file, "_");
+         else
+           fprintf (file, "%%r%d", regno);
+       }
     }
   else if (subreg_offset >= 0)
     {
...

[Bug target/104815] [nvptx] Use bitbucket operand when REG_UNUSED

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104815

--- Comment #1 from Tom de Vries <vries at gcc dot gnu.org> ---
With the tentative patch, I'm running into:
...
ptxas 20000224-1.o, line 72; error   : Result discard mode is not allowed for
instruction 'ld'
nvptx-as: ptxas terminated with signal 11 [Segmentation fault], core dumped
...

For:
...
        {
                .param.u32 %value_in;                                           
                .param.u64 %out_arg1;                                           
                st.param.u64 [%out_arg1], %r61;                                 
                call (%value_in), call_critical_lisp_code, (%out_arg1);         
                ld.param.u32    _, [%value_in];                                 
        }
...

Now trying a more restrictive approach, using an operand modifier x, only used
on atom operations, like so:
...
-      = "%.\\tatom%A1.add%t0\\t%0, %1, %2;";
+      = "%.\\tatom%A1.add%t0\\t%x0, %1, %2;";
...

[Bug target/104815] [nvptx] Use bitbucket operand when REG_UNUSED

[cvs-commit at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104815

--- Comment #2 from CVS Commits <cvs-commit at gcc dot gnu.org> ---
The master branch has been updated by Tom de Vries <vries@gcc.gnu.org>:

https://gcc.gnu.org/g:3ebcc053a4bd32973762b671b444730baf558805

commit r12-7584-g3ebcc053a4bd32973762b671b444730baf558805
Author: Tom de Vries <tdevries@suse.de>
Date:   Mon Mar 7 14:23:03 2022 +0100

    [nvptx] Use bit-bucket operand for atom insns

    For an atomic fetch operation that doesn't use the result:
    ...
      __atomic_fetch_add (p64, v64, MEMMODEL_RELAXED);
    ...
    we currently emit:
    ...
      atom.add.u64 %r26, [%r25], %r27;
    ...

    Detect the REG_UNUSED reg-note for %r26, and emit instead:
    ...
      atom.add.u64 _, [%r25], %r27;
    ...

    Likewise for all atom insns.

    Tested on nvptx.

    gcc/ChangeLog:

    2022-03-07  Tom de Vries  <tdevries@suse.de>

            PR target/104815
            * config/nvptx/nvptx.cc (nvptx_print_operand): Handle 'x' operand
            modifier.
            * config/nvptx/nvptx.md: Use %x0 destination operand in atom insns.

    gcc/testsuite/ChangeLog:

    2022-03-07  Tom de Vries  <tdevries@suse.de>

            PR target/104815
            * gcc.target/nvptx/atomic-bit-bucket-dest.c: New test.

[Bug target/104815] [nvptx] Use bitbucket operand when REG_UNUSED

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104815

Tom de Vries <vries at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
             Status|UNCONFIRMED                 |RESOLVED
   Target Milestone|---                         |12.0
         Resolution|---                         |FIXED

--- Comment #3 from Tom de Vries <vries at gcc dot gnu.org> ---
Fixed.