From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1643) id 5289B3858006; Wed, 14 Sep 2022 10:09:16 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 5289B3858006 DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1663150156; bh=R7RFuNtNIk9e1rtxw6cBXCNLv7WsHKOH9B6QrO8CjfQ=; h=From:To:Subject:Date:From; b=Yub+2CkI9XahMSs1HBkylJbuFSxn00h2tUIbBcBKIhDz0YwK3y/OmEnlZUD9eIBUg 4YJg+eyoOeoS75jmZSHCs3wak7x+D+VJT1tCl9Uc+W4U6ya81mB/Y6/xx+KVZ8udiq DnhLbj/mIoGSl15JC03IcEsxYmLR1AA+ZTkazLW8= MIME-Version: 1.0 Content-Transfer-Encoding: 7bit Content-Type: text/plain; charset="utf-8" From: Thomas Schwinge To: gcc-cvs@gcc.gnu.org Subject: [gcc r10-10986] [nvptx] Add bar.warp.sync X-Act-Checkin: gcc X-Git-Author: Tom de Vries X-Git-Refname: refs/heads/releases/gcc-10 X-Git-Oldrev: ddeca5944408edfac28b3d90b02b2d0b4cfe73d1 X-Git-Newrev: 50d326b90572d3e25a468f38603a6a971ecda938 Message-Id: <20220914100916.5289B3858006@sourceware.org> Date: Wed, 14 Sep 2022 10:09:16 +0000 (GMT) List-Id: https://gcc.gnu.org/g:50d326b90572d3e25a468f38603a6a971ecda938 commit r10-10986-g50d326b90572d3e25a468f38603a6a971ecda938 Author: Tom de Vries 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 * 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))]