[committed,nvptx] Add bar.warp.sync
Commit Message
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(+)
Comments
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
@@ -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),
@@ -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))]