[nvptx] Fix neutering of bb with only cond jump

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.

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

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@259125 138bc75d-0d04-0410-961f-82ee72b054a4

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 6e299cb33a759654eff98b1ba40ad55f83ef853e..517ac4e57ae6afb82b74638350f9770ee2322e91 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,9 @@
+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.
+
 2018-04-05  Shiva Chen  <shiva0217@gmail.com>
            Kito Cheng  <kito.cheng@gmail.com>
 

diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index b2b150f11d7cb8a6f5d1667f7af7a1a6e7286da7..a9a3053f3365d022da7bf9a636b11a5542214d80 100644 (file)
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -4048,6 +4048,7 @@ 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_insn *neuter_start = NULL;
   for (mode = GOMP_DIM_WORKER; mode <= GOMP_DIM_VECTOR; mode++)
     if (GOMP_DIM_MASK (mode) & skip_mask)
       {
@@ -4065,7 +4066,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)

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index a5a5e0631b223aff1f55dd33fc5fed2cf649c844..ea28859efda63473794966aac563af757805d33c 100644 (file)
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,8 @@
+2018-04-05  Tom de Vries  <tom@codesourcery.com>
+
+       PR target/85204
+       * testsuite/libgomp.oacc-c-c++-common/broadcast-1.c: New test.
+
 2018-03-26  Tom de Vries  <tom@codesourcery.com>
 
        PR tree-optimization/85063

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 (file)
index 0000000..ca0d37b
--- /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;
+}