[gomp4] propagating conditionals in worker-vector partitioned loops

[gomp4] propagating conditionals in worker-vector partitioned loops

[Cesar Philippidis]

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

I've applied the patch to gomp-4_0-branch to correct an issue involving
the propagation of variables used in conditional expressions to worker
and vector partitioned loops. More details regarding this patch can be
found here <https://gcc.gnu.org/ml/gcc-patches/2016-10/msg02187.html>

Cesar

[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: nvptx-broadcast.diff --]
[-- Type: text/x-patch; name="nvptx-broadcast.diff", Size: 2968 bytes --]

2016-10-26  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/
	* config/nvptx/nvptx.c (nvptx_single): Use a single predicate
	for loops partitioned across both worker and vector axes.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/broadcast-1.c: New test.


diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index 7bf5987..4e6ed60 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -3507,11 +3507,38 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
   /* Insert the vector test inside the worker test.  */
   unsigned mode;
   rtx_insn *before = tail;
+  rtx wvpred = NULL_RTX;
+  bool skip_vector = false;
+
+  /* Create a single predicate for loops containing both worker and
+     vectors.  */
+  if (cond_branch
+      && (GOMP_DIM_MASK (GOMP_DIM_WORKER) & mask)
+      && (GOMP_DIM_MASK (GOMP_DIM_VECTOR) & mask))
+    {
+      rtx regx = gen_reg_rtx (SImode);
+      rtx regy = gen_reg_rtx (SImode);
+      rtx tmp = gen_reg_rtx (SImode);
+      wvpred = gen_reg_rtx (BImode);
+
+      emit_insn_before (gen_oacc_dim_pos (regx, const1_rtx), head);
+      emit_insn_before (gen_oacc_dim_pos (regy, const2_rtx), head);
+      emit_insn_before (gen_rtx_SET (tmp, gen_rtx_IOR (SImode, regx, regy)),
+			head);
+      emit_insn_before (gen_rtx_SET (wvpred, gen_rtx_NE (BImode, tmp,
+							 const0_rtx)),
+			head);
+
+      skip_mask &= ~(GOMP_DIM_MASK (GOMP_DIM_VECTOR));
+      skip_vector = true;
+    }
+
   for (mode = GOMP_DIM_WORKER; mode <= GOMP_DIM_VECTOR; mode++)
     if (GOMP_DIM_MASK (mode) & skip_mask)
       {
 	rtx_code_label *label = gen_label_rtx ();
-	rtx pred = cfun->machine->axis_predicate[mode - GOMP_DIM_WORKER];
+	rtx pred = skip_vector ? wvpred
+	  : cfun->machine->axis_predicate[mode - GOMP_DIM_WORKER];
 
 	if (!pred)
 	  {
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/broadcast-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/broadcast-1.c
new file mode 100644
index 0000000..4dcb60d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/broadcast-1.c
@@ -0,0 +1,49 @@
+/* Ensure that worker-vector state conditional expressions are
+   properly handled by the nvptx backend.  */
+
+#include <assert.h>
+#include <math.h>
+
+
+#define N 1024
+
+int A[N][N] ;
+
+void test(int x)
+{
+#pragma acc parallel  num_gangs(16) num_workers(4) vector_length(32) copyout(A)
+  {
+#pragma acc loop gang
+    for(int j=0;j<N;j++)
+      {
+	if (x==1)
+	  {
+#pragma acc loop worker vector
+	    for(int i=0;i<N;i++)
+	      A[i][j] = 1;
+	  }
+	else
+	  {
+#pragma acc loop worker vector
+	    for(int i=0;i<N;i++)
+	      A[i][j] = -1;
+	  }
+      }
+  }
+}
+
+
+int main(void)
+{
+  test (0);
+  for (int i = 0; i < N; i++)
+    for (int j = 0; j < N; j++)
+      assert (A[i][j] == -1);
+
+  test (1);
+  for (int i = 0; i < N; i++)
+    for (int j = 0; j < N; j++)
+      assert (A[i][j] == 1);
+
+  return 0;
+}

[og7] Backport "[nvptx] Fix neutering of bb with only cond jump"

[Tom de Vries]

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

[ was: Re: [gomp4] propagating conditionals in worker-vector partitioned 
loops ]

On 10/28/2016 06:33 PM, Cesar Philippidis wrote:
> I've applied the patch to gomp-4_0-branch to correct an issue involving
> the propagation of variables used in conditional expressions to worker
> and vector partitioned loops. More details regarding this patch can be
> found here<https://gcc.gnu.org/ml/gcc-patches/2016-10/msg02187.html>

I've reverted this patch on og7, and backported the fix for PR85204.

Thanks,
- Tom

[-- Attachment #2: 0001-Backport-nvptx-Fix-neutering-of-bb-with-only-cond-jump.patch --]
[-- Type: text/x-patch, Size: 2556 bytes --]

Backport "[nvptx] Fix neutering of bb with only cond jump"

2018-04-11  Tom de Vries  <tom@codesourcery.com>

	backport from trunk:
	2018-04-05  Tom de Vries  <tom@codesourcery.com>

	PR target/85204
	* config/nvptx/nvptx.c (nvptx_single): Fix neutering of bb with only
	cond jump.

	revert:
	2016-10-28  Cesar Philippidis  <cesar@codesourcery.com>

	* config/nvptx/nvptx.c (nvptx_single): Use a single predicate
	for loops partitioned across both worker and vector axes.

---
 gcc/config/nvptx/nvptx.c | 35 ++++++-----------------------------
 1 file changed, 6 insertions(+), 29 deletions(-)

diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index cd89d17..3c48c14 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -4196,38 +4196,12 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
   /* Insert the vector test inside the worker test.  */
   unsigned mode;
   rtx_insn *before = tail;
-  rtx wvpred = NULL_RTX;
-  bool skip_vector = false;
-
-  /* Create a single predicate for loops containing both worker and
-     vectors.  */
-  if (cond_branch
-      && (GOMP_DIM_MASK (GOMP_DIM_WORKER) & mask)
-      && (GOMP_DIM_MASK (GOMP_DIM_VECTOR) & mask))
-    {
-      rtx regx = gen_reg_rtx (SImode);
-      rtx regy = gen_reg_rtx (SImode);
-      rtx tmp = gen_reg_rtx (SImode);
-      wvpred = gen_reg_rtx (BImode);
-
-      emit_insn_before (gen_oacc_dim_pos (regx, const1_rtx), head);
-      emit_insn_before (gen_oacc_dim_pos (regy, const2_rtx), head);
-      emit_insn_before (gen_rtx_SET (tmp, gen_rtx_IOR (SImode, regx, regy)),
-			head);
-      emit_insn_before (gen_rtx_SET (wvpred, gen_rtx_NE (BImode, tmp,
-							 const0_rtx)),
-			head);
-
-      skip_mask &= ~(GOMP_DIM_MASK (GOMP_DIM_VECTOR));
-      skip_vector = true;
-    }
-
+  rtx_insn *neuter_start = NULL;
   for (mode = GOMP_DIM_WORKER; mode <= GOMP_DIM_VECTOR; mode++)
     if (GOMP_DIM_MASK (mode) & skip_mask)
       {
 	rtx_code_label *label = gen_label_rtx ();
-	rtx pred = skip_vector ? wvpred
-	  : cfun->machine->axis_predicate[mode - GOMP_DIM_WORKER];
+	rtx pred = cfun->machine->axis_predicate[mode - GOMP_DIM_WORKER];
 
 	if (!pred)
 	  {
@@ -4240,7 +4214,10 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
 	  br = gen_br_true (pred, label);
 	else
 	  br = gen_br_true_uni (pred, label);
-	emit_insn_before (br, head);
+	if (neuter_start)
+	  neuter_start = emit_insn_after (br, neuter_start);
+	else
+	  neuter_start = emit_insn_before (br, head);
 
 	LABEL_NUSES (label)++;
 	if (tail_branch)