public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc r10-10986] [nvptx] Add bar.warp.sync
@ 2022-09-14 10:09 Thomas Schwinge
0 siblings, 0 replies; only message in thread
From: Thomas Schwinge @ 2022-09-14 10:09 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:50d326b90572d3e25a468f38603a6a971ecda938
commit r10-10986-g50d326b90572d3e25a468f38603a6a971ecda938
Author: Tom de Vries <tdevries@suse.de>
Date: Thu Jan 27 15:03:59 2022 +0100
[nvptx] Add bar.warp.sync
On a GT 1030 (sm_61), with driver version 470.94 I run into:
...
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims.c \
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none \
-O2 execution test
...
which minimizes to the same test-case as listed in commit "[nvptx] Update
default ptx isa to 6.3".
The first divergent branch looks like:
...
{
.reg .u32 %x;
mov.u32 %x,%tid.x;
setp.ne.u32 %r59,%x,0;
}
@ %r59 bra $L15;
mov.u64 %r48,%ar0;
mov.u32 %r22,2;
ld.u64 %r53,[%r48];
mov.u32 %r55,%r22;
mov.u32 %r54,1;
$L15:
...
and when inspecting the generated SASS, the branch is not setup as a divergent
branch, but instead as a regular branch.
This causes us to execute a shfl.sync insn in divergent mode, which is likely
to cause trouble given a remark in the ptx isa version 6.3, which mentions
that for .target sm_6x or below, all threads must excute the same
shfl.sync instruction in convergence.
Fix this by placing a "bar.warp.sync 0xffffffff" at the desired convergence
point (in the example above, after $L15).
Tested on x86_64 with nvptx accelerator.
gcc/ChangeLog:
2022-01-31 Tom de Vries <tdevries@suse.de>
* config/nvptx/nvptx.c (nvptx_single): Use nvptx_warpsync.
* config/nvptx/nvptx.md (define_c_enum "unspecv"): Add
UNSPECV_WARPSYNC.
(define_insn "nvptx_warpsync"): New define_insn.
(cherry picked from commit bba61d403d05202deb698b352a4faef3feb1f04d)
Diff:
---
gcc/config/nvptx/nvptx.c | 7 +++++++
gcc/config/nvptx/nvptx.md | 7 +++++++
2 files changed, 14 insertions(+)
diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index 243bf9da51a..eec07fc3d13 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -4392,6 +4392,7 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
rtx_insn *neuter_start = NULL;
rtx_insn *worker_label = NULL, *vector_label = NULL;
rtx_insn *worker_jump = NULL, *vector_jump = NULL;
+ rtx_insn *warp_sync = NULL;
for (mode = GOMP_DIM_WORKER; mode <= GOMP_DIM_VECTOR; mode++)
if (GOMP_DIM_MASK (mode) & skip_mask)
{
@@ -4424,11 +4425,15 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
if (tail_branch)
{
label_insn = emit_label_before (label, before);
+ if (TARGET_PTX_6_0 && mode == GOMP_DIM_VECTOR)
+ warp_sync = emit_insn_after (gen_nvptx_warpsync (), label_insn);
before = label_insn;
}
else
{
label_insn = emit_label_after (label, tail);
+ if (TARGET_PTX_6_0 && mode == GOMP_DIM_VECTOR)
+ warp_sync = emit_insn_after (gen_nvptx_warpsync (), label_insn);
if ((mode == GOMP_DIM_VECTOR || mode == GOMP_DIM_WORKER)
&& CALL_P (tail) && find_reg_note (tail, REG_NORETURN, NULL))
emit_insn_after (gen_exit (), label_insn);
@@ -4496,6 +4501,8 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
setp.ne.u32 %rcond,%rcondu32,0;
*/
rtx_insn *label = PREV_INSN (tail);
+ if (label == warp_sync)
+ label = PREV_INSN (label);
gcc_assert (label && LABEL_P (label));
rtx tmp = gen_reg_rtx (BImode);
emit_insn_before (gen_movbi (tmp, const0_rtx),
diff --git a/gcc/config/nvptx/nvptx.md b/gcc/config/nvptx/nvptx.md
index 089cdf04f2e..b2350ab0d2d 100644
--- a/gcc/config/nvptx/nvptx.md
+++ b/gcc/config/nvptx/nvptx.md
@@ -55,6 +55,7 @@
UNSPECV_CAS
UNSPECV_XCHG
UNSPECV_BARSYNC
+ UNSPECV_WARPSYNC
UNSPECV_MEMBAR
UNSPECV_MEMBAR_CTA
UNSPECV_DIM_POS
@@ -1476,6 +1477,12 @@
}
[(set_attr "predicable" "false")])
+(define_insn "nvptx_warpsync"
+ [(unspec_volatile [(const_int 0)] UNSPECV_WARPSYNC)]
+ "TARGET_PTX_6_0"
+ "\\tbar.warp.sync\\t0xffffffff;"
+ [(set_attr "predicable" "false")])
+
(define_expand "memory_barrier"
[(set (match_dup 0)
(unspec_volatile:BLK [(match_dup 0)] UNSPECV_MEMBAR))]
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2022-09-14 10:09 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-09-14 10:09 [gcc r10-10986] [nvptx] Add bar.warp.sync Thomas Schwinge
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).