From mboxrd@z Thu Jan  1 00:00:00 1970
Return-Path: <gcc-bugzilla@gcc.gnu.org>
Received: by sourceware.org (Postfix, from userid 48)
 id AB4403858428; Mon,  7 Mar 2022 10:16:15 +0000 (GMT)
DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org AB4403858428
From: "vries at gcc dot gnu.org" <gcc-bugzilla@gcc.gnu.org>
To: gcc-bugs@gcc.gnu.org
Subject: [Bug target/104815] New: [nvptx] Use bitbucket operand when REG_UNUSED
Date: Mon, 07 Mar 2022 10:16:15 +0000
X-Bugzilla-Reason: CC
X-Bugzilla-Type: new
X-Bugzilla-Watch-Reason: None
X-Bugzilla-Product: gcc
X-Bugzilla-Component: target
X-Bugzilla-Version: 12.0
X-Bugzilla-Keywords: 
X-Bugzilla-Severity: enhancement
X-Bugzilla-Who: vries at gcc dot gnu.org
X-Bugzilla-Status: UNCONFIRMED
X-Bugzilla-Resolution: 
X-Bugzilla-Priority: P3
X-Bugzilla-Assigned-To: unassigned at gcc dot gnu.org
X-Bugzilla-Target-Milestone: ---
X-Bugzilla-Flags: 
X-Bugzilla-Changed-Fields: bug_id short_desc product version bug_status
 bug_severity priority component assigned_to reporter target_milestone
Message-ID: <bug-104815-4@http.gcc.gnu.org/bugzilla/>
Content-Type: text/plain; charset="UTF-8"
Content-Transfer-Encoding: quoted-printable
X-Bugzilla-URL: http://gcc.gnu.org/bugzilla/
Auto-Submitted: auto-generated
MIME-Version: 1.0
X-BeenThere: gcc-bugs@gcc.gnu.org
X-Mailman-Version: 2.1.29
Precedence: list
List-Id: Gcc-bugs mailing list <gcc-bugs.gcc.gnu.org>
List-Unsubscribe: <https://gcc.gnu.org/mailman/options/gcc-bugs>,
 <mailto:gcc-bugs-request@gcc.gnu.org?subject=unsubscribe>
List-Archive: <https://gcc.gnu.org/pipermail/gcc-bugs/>
List-Post: <mailto:gcc-bugs@gcc.gnu.org>
List-Help: <mailto:gcc-bugs-request@gcc.gnu.org?subject=help>
List-Subscribe: <https://gcc.gnu.org/mailman/listinfo/gcc-bugs>,
 <mailto:gcc-bugs-request@gcc.gnu.org?subject=subscribe>
X-List-Received-Date: Mon, 07 Mar 2022 10:16:15 -0000

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

            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 =3D 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 =3D regno_reg_rtx[regno];
+         if (current_output_insn
+             && (find_reg_note (current_output_insn, REG_UNUSED, reg)
+                 !=3D NULL_RTX))
+           fprintf (file, "_");
+         else
+           fprintf (file, "%%r%d", regno);
+       }
     }
   else if (subreg_offset >=3D 0)
     {
...=