| Message ID | 20220201183111.GA4257@delia.home |
|---|---|
| State | New |
| Headers | show |
| Series | [committed,nvptx] Add bar.warp.sync | expand |
Hi Tom! On 2022-02-01T19:31:13+0100, Tom de Vries via Gcc-patches <gcc-patches@gcc.gnu.org> wrote: > 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 > ... This relates to PR99932 "OpenACC/nvptx offloading execution regressions starting with CUDA 11.2-era Nvidia Driver 460.27.04". You've fixed that for GCC 12+, but testing nvptx offloading for GCC 11, GCC 10 branches on a system with current Nvidia Driver, we're still running into the several FAILs and annoying 'WARNING: program timed out' reported in PR99932. Given that GCC 11, GCC 10 only build '.version 3.1' target libraries, the patch below is actually a no-op ('!TARGET_PTX_6_0'; thus the attached "nvptx: Define (dummy) 'TARGET_PTX_6_0'"), but having it makes it easier to cherry-pick the actual relevant "[nvptx] Add uniform_warp_check insn". OK to push to GCC 11, GCC 10 branches the attached "nvptx: Define (dummy) 'TARGET_PTX_6_0'", "[nvptx] Add bar.warp.sync"? Grüße Thomas > 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. > > Committed to trunk. > > Thanks, > - Tom > > [nvptx] Add bar.warp.sync > > gcc/ChangeLog: > > 2022-01-31 Tom de Vries <tdevries@suse.de> > > * config/nvptx/nvptx.cc (nvptx_single): Use nvptx_warpsync. > * config/nvptx/nvptx.md (define_c_enum "unspecv"): Add > UNSPECV_WARPSYNC. > (define_insn "nvptx_warpsync"): New define_insn. > > --- > gcc/config/nvptx/nvptx.cc | 7 +++++++ > gcc/config/nvptx/nvptx.md | 7 +++++++ > 2 files changed, 14 insertions(+) > > diff --git a/gcc/config/nvptx/nvptx.cc b/gcc/config/nvptx/nvptx.cc > index 262e8f9cc1b..1b91990ca1f 100644 > --- a/gcc/config/nvptx/nvptx.cc > +++ b/gcc/config/nvptx/nvptx.cc > @@ -4598,6 +4598,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) > { > @@ -4630,11 +4631,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); > @@ -4702,6 +4707,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 b39116520ba..b4c7cd6e56d 100644 > --- a/gcc/config/nvptx/nvptx.md > +++ b/gcc/config/nvptx/nvptx.md > @@ -56,6 +56,7 @@ (define_c_enum "unspecv" [ > UNSPECV_CAS > UNSPECV_XCHG > UNSPECV_BARSYNC > + UNSPECV_WARPSYNC > UNSPECV_MEMBAR > UNSPECV_MEMBAR_CTA > UNSPECV_MEMBAR_GL > @@ -1978,6 +1979,12 @@ (define_insn "nvptx_barsync" > } > [(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))] ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
On 9/14/22 11:41, Thomas Schwinge wrote: > Hi Tom! > > On 2022-02-01T19:31:13+0100, Tom de Vries via Gcc-patches <gcc-patches@gcc.gnu.org> wrote: >> 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 >> ... > > This relates to PR99932 "OpenACC/nvptx offloading execution regressions > starting with CUDA 11.2-era Nvidia Driver 460.27.04". You've fixed that > for GCC 12+, but testing nvptx offloading for GCC 11, GCC 10 branches on > a system with current Nvidia Driver, we're still running into the several > FAILs and annoying 'WARNING: program timed out' reported in PR99932. > > Given that GCC 11, GCC 10 only build '.version 3.1' target libraries, the > patch below is actually a no-op ('!TARGET_PTX_6_0'; thus the attached > "nvptx: Define (dummy) 'TARGET_PTX_6_0'"), but having it makes it easier > to cherry-pick the actual relevant "[nvptx] Add uniform_warp_check insn". > > OK to push to GCC 11, GCC 10 branches the attached > "nvptx: Define (dummy) 'TARGET_PTX_6_0'", "[nvptx] Add bar.warp.sync"? > LGTM, and thanks for taking care of this. Thanks, - Tom
diff --git a/gcc/config/nvptx/nvptx.cc b/gcc/config/nvptx/nvptx.cc index 262e8f9cc1b..1b91990ca1f 100644 --- a/gcc/config/nvptx/nvptx.cc +++ b/gcc/config/nvptx/nvptx.cc @@ -4598,6 +4598,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) { @@ -4630,11 +4631,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); @@ -4702,6 +4707,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 b39116520ba..b4c7cd6e56d 100644 --- a/gcc/config/nvptx/nvptx.md +++ b/gcc/config/nvptx/nvptx.md @@ -56,6 +56,7 @@ (define_c_enum "unspecv" [ UNSPECV_CAS UNSPECV_XCHG UNSPECV_BARSYNC + UNSPECV_WARPSYNC UNSPECV_MEMBAR UNSPECV_MEMBAR_CTA UNSPECV_MEMBAR_GL @@ -1978,6 +1979,12 @@ (define_insn "nvptx_barsync" } [(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))]
Hi, 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. Committed to trunk. Thanks, - Tom [nvptx] Add bar.warp.sync gcc/ChangeLog: 2022-01-31 Tom de Vries <tdevries@suse.de> * config/nvptx/nvptx.cc (nvptx_single): Use nvptx_warpsync. * config/nvptx/nvptx.md (define_c_enum "unspecv"): Add UNSPECV_WARPSYNC. (define_insn "nvptx_warpsync"): New define_insn. --- gcc/config/nvptx/nvptx.cc | 7 +++++++ gcc/config/nvptx/nvptx.md | 7 +++++++ 2 files changed, 14 insertions(+)