public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* [committed][nvptx] Add bar.warp.sync
@ 2022-02-01 18:31 Tom de Vries
  2022-09-14  9:41 ` Thomas Schwinge
  0 siblings, 1 reply; 3+ messages in thread
From: Tom de Vries @ 2022-02-01 18:31 UTC (permalink / raw)
  To: gcc-patches

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(+)

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))]

^ permalink raw reply	[flat|nested] 3+ messages in thread

* Re: [committed][nvptx] Add bar.warp.sync
  2022-02-01 18:31 [committed][nvptx] Add bar.warp.sync Tom de Vries
@ 2022-09-14  9:41 ` Thomas Schwinge
  2022-09-14  9:54   ` Tom de Vries
  0 siblings, 1 reply; 3+ messages in thread
From: Thomas Schwinge @ 2022-09-14  9:41 UTC (permalink / raw)
  To: Tom de Vries; +Cc: gcc-patches

[-- Attachment #1: Type: text/plain, Size: 5562 bytes --]

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

[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-nvptx-Define-dummy-TARGET_PTX_6_0.patch --]
[-- Type: text/x-diff, Size: 996 bytes --]

From 03b69e5c334e8bb17306961e81b4410ba7408774 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@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


[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #3: 0002-nvptx-Add-bar.warp.sync.patch --]
[-- Type: text/x-diff, Size: 4042 bytes --]

From 3951a2a1c10f6d583d18b2c965165cc1f185125e Mon Sep 17 00:00:00 2001
From: Tom de Vries <tdevries@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  <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)
---
 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


^ permalink raw reply	[flat|nested] 3+ messages in thread

* Re: [committed][nvptx] Add bar.warp.sync
  2022-09-14  9:41 ` Thomas Schwinge
@ 2022-09-14  9:54   ` Tom de Vries
  0 siblings, 0 replies; 3+ messages in thread
From: Tom de Vries @ 2022-09-14  9:54 UTC (permalink / raw)
  To: Thomas Schwinge; +Cc: gcc-patches

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


^ permalink raw reply	[flat|nested] 3+ messages in thread

end of thread, other threads:[~2022-09-14  9:54 UTC | newest]

Thread overview: 3+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-02-01 18:31 [committed][nvptx] Add bar.warp.sync Tom de Vries
2022-09-14  9:41 ` Thomas Schwinge
2022-09-14  9:54   ` Tom de Vries

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).