diff mbox

[gomp4,nvptx,committed] Fix assert in nvptx_propagate_unified

Message ID b3744724-1127-a5ec-5502-c9076f0b3d31@mentor.com
State New
Headers show

Commit Message

Tom de Vries June 30, 2017, 3:15 p.m. UTC 
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
diff mbox

Patch

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;
+}