[gomp4, nvptx, committed] Fix assert in nvptx_propagate_unified

[Tom de Vries <Tom_deVries at mentor dot com>]

Hi,

with the openacc test-case in attached patch, I ran into an assert here:
...
static void
nvptx_propagate_unified (rtx_insn *unified)
{
   rtx_insn *probe = unified;
   rtx cond_reg = SET_DEST (PATTERN (unified));
   rtx pat;

   /* Find the comparison.  (We could skip this and simply scan to he
      blocks' terminating branch, if we didn't care for self
      checking.)  */
   for (;;)
     {
       probe = NEXT_INSN (probe);
       pat = PATTERN (probe);

       if (GET_CODE (pat) == SET
           && GET_RTX_CLASS (GET_CODE (SET_SRC (pat))) == RTX_COMPARE
           && XEXP (SET_SRC (pat), 0) == cond_reg)
     break;
       gcc_assert (NONJUMP_INSN_P (probe));
     }
...

The assert happens when processing insn 56:
...
(insn 54 53 56 3 (set (reg:SI 47 [ _71 ])
          (unspec:SI [
                  (reg:SI 36 [ _58 ])
              ] UNSPEC_BR_UNIFIED)) 108 {cond_uni}
       (nil))
(note 56 54 57 3 NOTE_INSN_DELETED)
(insn 57 56 58 3 (set (reg:BI 68)
          (gt:BI (reg:SI 47 [ _71 ])
              (const_int 1 [0x1]))) 99 {*cmpsi}
       (expr_list:REG_DEAD (reg:SI 47 [ _71 ])
          (nil)))
...
The insn 56 was originally a '(set (reg x) (const_int 1))', but that one 
has been combined into insn 57 and replaced with a 'NOTE_INSN_DELETED'. 
So it seems reasonable for the loop to skip over this note.

Fixed by making the assert condition less strict.

Build on x86_64 with nvptx accelerator.

Tested test-case included in the patch.

Committed as trivial.

Thanks,
- Tom

Fix assert in nvptx_propagate_unified

2017-06-30  Tom de Vries  <tom@codesourcery.com>

	* config/nvptx/nvptx.c (nvptx_propagate_unified): Fix gcc_assert
	condition by allowing !INSN_P.

	* testsuite/libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c: New test.

---
 gcc/config/nvptx/nvptx.c                           |  2 +-
 .../reduction-cplx-flt-2.c                         | 32 ++++++++++++++++++++++
 2 files changed, 33 insertions(+), 1 deletion(-)

diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index 4a93bba..a3abb4b 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -2300,7 +2300,7 @@ nvptx_propagate_unified (rtx_insn *unified)
 	  && GET_RTX_CLASS (GET_CODE (SET_SRC (pat))) == RTX_COMPARE
 	  && XEXP (SET_SRC (pat), 0) == cond_reg)
 	break;
-      gcc_assert (NONJUMP_INSN_P (probe));
+      gcc_assert (NONJUMP_INSN_P (probe) || !INSN_P (probe));
     }
   rtx pred_reg = SET_DEST (pat);
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c
new file mode 100644
index 0000000..f80f38c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c
@@ -0,0 +1,32 @@
+#include <complex.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+typedef float _Complex Type;
+
+#define N 32
+
+int
+main (void)
+{
+  Type ary[N];
+
+  for (int ix = 0; ix < N;  ix++)
+    ary[ix] = 1.0 + 1.0i;
+
+  Type tprod = 1.0;
+
+#pragma acc parallel vector_length(32)
+  {
+#pragma acc loop vector reduction (*:tprod)
+    for (int ix = 0; ix < N; ix++)
+      tprod *= ary[ix];
+  }
+
+  Type expected = 65536.0;
+
+  if (tprod != expected)
+    abort ();
+
+  return 0;
+}