[nvptx] Handle unused result in nvptx_unisimt_handle_set
For an example: ... #pragma omp target map(tofrom: counter_N0) #pragma omp simd for (int i = 0 ; i < 1 ; i++ ) { #pragma omp atomic update counter_N0 = counter_N0 + 1 ; } ... I noticed that the result of the atomic update (%r30) is propagated: ... @%r33 atom.add.u32 _, [%r29], 1; shfl.sync.idx.b32 %r30, %r30, %r32, 31, 0xffffffff; ... even though it is unused (which is why the bit bucket operand _ is used). Fix this by not emitting the shuffle in this case, such that we have instead: ... @%r33 atom.add.u32 _, [%r29], 1; bar.warp.sync 0xffffffff; ... Tested on nvptx. gcc/ChangeLog: 2022-03-07 Tom de Vries <tdevries@suse.de> * config/nvptx/nvptx.cc (nvptx_unisimt_handle_set): Handle unused result. gcc/testsuite/ChangeLog: 2022-03-07 Tom de Vries <tdevries@suse.de> * gcc.target/nvptx/uniform-simt-4.c: New test.
This commit is contained in:
parent
3ebcc053a4
commit
3e743d654b
@ -3274,7 +3274,9 @@ static bool
|
||||
nvptx_unisimt_handle_set (rtx set, rtx_insn *insn, rtx master)
|
||||
{
|
||||
rtx reg;
|
||||
if (GET_CODE (set) == SET && REG_P (reg = SET_DEST (set)))
|
||||
if (GET_CODE (set) == SET
|
||||
&& REG_P (reg = SET_DEST (set))
|
||||
&& find_reg_note (insn, REG_UNUSED, reg) == NULL_RTX)
|
||||
{
|
||||
emit_insn_after (nvptx_gen_shuffle (reg, reg, master, SHUFFLE_IDX),
|
||||
insn);
|
||||
|
22
gcc/testsuite/gcc.target/nvptx/uniform-simt-4.c
Normal file
22
gcc/testsuite/gcc.target/nvptx/uniform-simt-4.c
Normal file
@ -0,0 +1,22 @@
|
||||
/* { dg-do compile } */
|
||||
/* { dg-options "-O2 -muniform-simt -mptx=_" } */
|
||||
|
||||
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;
|
||||
}
|
||||
|
||||
/* { dg-final { scan-assembler-times "atom.add.u64\[\t \]+_," 1 } } */
|
||||
/* { dg-final { scan-assembler-times "bar.warp.sync" 1 } } */
|
||||
/* { dg-final { scan-assembler-not "shfl.sync.idx" } } */
|
Loading…
Reference in New Issue
Block a user