public inbox for gcc-bugs@sourceware.org
help / color / mirror / Atom feed
* [Bug target/104815] New: [nvptx] Use bitbucket operand when REG_UNUSED
@ 2022-03-07 10:16 vries at gcc dot gnu.org
2022-03-07 14:34 ` [Bug target/104815] " vries at gcc dot gnu.org
` (2 more replies)
0 siblings, 3 replies; 4+ messages in thread
From: vries at gcc dot gnu.org @ 2022-03-07 10:16 UTC (permalink / raw)
To: gcc-bugs
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)
{
...
^ permalink raw reply [flat|nested] 4+ messages in thread
* [Bug target/104815] [nvptx] Use bitbucket operand when REG_UNUSED
2022-03-07 10:16 [Bug target/104815] New: [nvptx] Use bitbucket operand when REG_UNUSED vries at gcc dot gnu.org
@ 2022-03-07 14:34 ` vries at gcc dot gnu.org
2022-03-10 11:22 ` cvs-commit at gcc dot gnu.org
2022-03-10 11:30 ` vries at gcc dot gnu.org
2 siblings, 0 replies; 4+ messages in thread
From: vries at gcc dot gnu.org @ 2022-03-07 14:34 UTC (permalink / raw)
To: gcc-bugs
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;";
...
^ permalink raw reply [flat|nested] 4+ messages in thread
* [Bug target/104815] [nvptx] Use bitbucket operand when REG_UNUSED
2022-03-07 10:16 [Bug target/104815] New: [nvptx] Use bitbucket operand when REG_UNUSED vries at gcc dot gnu.org
2022-03-07 14:34 ` [Bug target/104815] " vries at gcc dot gnu.org
@ 2022-03-10 11:22 ` cvs-commit at gcc dot gnu.org
2022-03-10 11:30 ` vries at gcc dot gnu.org
2 siblings, 0 replies; 4+ messages in thread
From: cvs-commit at gcc dot gnu.org @ 2022-03-10 11:22 UTC (permalink / raw)
To: gcc-bugs
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.
^ permalink raw reply [flat|nested] 4+ messages in thread
* [Bug target/104815] [nvptx] Use bitbucket operand when REG_UNUSED
2022-03-07 10:16 [Bug target/104815] New: [nvptx] Use bitbucket operand when REG_UNUSED vries at gcc dot gnu.org
2022-03-07 14:34 ` [Bug target/104815] " vries at gcc dot gnu.org
2022-03-10 11:22 ` cvs-commit at gcc dot gnu.org
@ 2022-03-10 11:30 ` vries at gcc dot gnu.org
2 siblings, 0 replies; 4+ messages in thread
From: vries at gcc dot gnu.org @ 2022-03-10 11:30 UTC (permalink / raw)
To: gcc-bugs
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.
^ permalink raw reply [flat|nested] 4+ messages in thread
end of thread, other threads:[~2022-03-10 11:30 UTC | newest]
Thread overview: 4+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-03-07 10:16 [Bug target/104815] New: [nvptx] Use bitbucket operand when REG_UNUSED vries at gcc dot gnu.org
2022-03-07 14:34 ` [Bug target/104815] " vries at gcc dot gnu.org
2022-03-10 11:22 ` cvs-commit at gcc dot gnu.org
2022-03-10 11:30 ` vries at gcc dot gnu.org
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).