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)
{
...
