public inbox for gcc-cvs@sourceware.org help / color / mirror / Atom feed
From: Thomas Schwinge <tschwinge@gcc.gnu.org> To: gcc-cvs@gcc.gnu.org Subject: [gcc r11-10256] [nvptx] Add bar.warp.sync Date: Wed, 14 Sep 2022 10:09:01 +0000 (GMT) [thread overview] Message-ID: <20220914100901.894D73858297@sourceware.org> (raw) https://gcc.gnu.org/g:8b447b4be099ed5854666d53e6aa555b225403ac commit r11-10256-g8b447b4be099ed5854666d53e6aa555b225403ac 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 7a7a9130e84..06861f6f405 100644 --- a/gcc/config/nvptx/nvptx.c +++ b/gcc/config/nvptx/nvptx.c @@ -4503,6 +4503,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) { @@ -4535,11 +4536,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); @@ -4607,6 +4612,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 0f15609ee4b..c33b5d3279f 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 @@ -1718,6 +1719,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))]
reply other threads:[~2022-09-14 10:09 UTC|newest] Thread overview: [no followups] expand[flat|nested] mbox.gz Atom feed
Reply instructions: You may reply publicly to this message via plain-text email using any one of the following methods: * Save the following mbox file, import it into your mail client, and reply-to-all from there: mbox Avoid top-posting and favor interleaved quoting: https://en.wikipedia.org/wiki/Posting_style#Interleaved_style * Reply using the --to, --cc, and --in-reply-to switches of git-send-email(1): git send-email \ --in-reply-to=20220914100901.894D73858297@sourceware.org \ --to=tschwinge@gcc.gnu.org \ --cc=gcc-cvs@gcc.gnu.org \ /path/to/YOUR_REPLY https://kernel.org/pub/software/scm/git/docs/git-send-email.html * If your mail client supports setting the In-Reply-To header via mailto: links, try the mailto: linkBe sure your reply has a Subject: header at the top and a blank line before the message body.
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).