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 <tdevr...@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
>From 03b69e5c334e8bb17306961e81b4410ba7408774 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <tho...@codesourcery.com> Date: Tue, 13 Sep 2022 17:26:33 +0200 Subject: [PATCH 1/2] nvptx: Define (dummy) 'TARGET_PTX_6_0' 'TARGET_PTX_*' not applicable before GCC 12, but makes it easier to cherry-pick patches. gcc/ * config/nvptx/nvptx.h (TARGET_PTX_6_0): '#define' to 'false'. --- gcc/config/nvptx/nvptx.h | 3 +++ 1 file changed, 3 insertions(+) diff --git a/gcc/config/nvptx/nvptx.h b/gcc/config/nvptx/nvptx.h index 6ebcc760771..ca1aa11d852 100644 --- a/gcc/config/nvptx/nvptx.h +++ b/gcc/config/nvptx/nvptx.h @@ -95,6 +95,9 @@ #define TARGET_SM35 (ptx_isa_option >= PTX_ISA_SM35) +/* 'TARGET_PTX_*' not applicable before GCC 12. */ +#define TARGET_PTX_6_0 false + /* Registers. Since ptx is a virtual target, we just define a few hard registers for special purposes and leave pseudos unallocated. We have to have some available hard registers, to keep gcc setup -- 2.35.1
>From 3951a2a1c10f6d583d18b2c965165cc1f185125e Mon Sep 17 00:00:00 2001 From: Tom de Vries <tdevr...@suse.de> Date: Thu, 27 Jan 2022 15:03:59 +0100 Subject: [PATCH 2/2] [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 <tdevr...@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) --- 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 e3e84dfd4e4..6782f13f787 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))] -- 2.35.1