* [gomp4, nvptx, committed] Fix assert in nvptx_propagate_unified
@ 2017-06-30 15:15 ` Tom de Vries
2017-07-06 7:16 ` Thomas Schwinge
0 siblings, 1 reply; 3+ messages in thread
From: Tom de Vries @ 2017-06-30 15:15 UTC (permalink / raw)
To: GCC Patches; +Cc: Thomas Schwinge
[-- Attachment #1: Type: text/plain, Size: 1532 bytes --]
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
[-- Attachment #2: 0001-Fix-assert-in-nvptx_propagate_unified.patch --]
[-- Type: text/x-patch, Size: 1730 bytes --]
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;
+}
^ permalink raw reply [flat|nested] 3+ messages in thread
* Re: [gomp4, nvptx, committed] Fix assert in nvptx_propagate_unified
2017-06-30 15:15 ` [gomp4, nvptx, committed] Fix assert in nvptx_propagate_unified Tom de Vries
@ 2017-07-06 7:16 ` Thomas Schwinge
2017-07-16 18:59 ` [gomp4, nvptx, committed] Fix --enable-checking=rtl " Tom de Vries
0 siblings, 1 reply; 3+ messages in thread
From: Thomas Schwinge @ 2017-07-06 7:16 UTC (permalink / raw)
To: Tom de Vries, GCC Patches
Hi Tom!
On Fri, 30 Jun 2017 17:15:24 +0200, Tom de Vries <Tom_deVries@mentor.com> wrote:
> with the openacc test-case in attached patch, I ran into an assert here:
Using your test case, in my build with
"--enable-checking=yes,df,fold,rtl", I already earlier run into an ICE...
> 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);
... here:
[...]/libgomp/testsuite/libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c:19:9: internal compiler error: RTL check: expected elt 3 type 'e' or 'u', have '0' (rtx note) in PATTERN, at rtl.h:1440
Breakpoint 2, internal_error (gmsgid=0x10cc840 "RTL check: expected elt %d type '%c' or '%c', have '%c' (rtx %s) in %s, at %s:%d") at [...]/gcc/diagnostic.c:1251
1251 {
(gdb) bt
#0 internal_error (gmsgid=0x10cc840 "RTL check: expected elt %d type '%c' or '%c', have '%c' (rtx %s) in %s, at %s:%d") at [...]/gcc/diagnostic.c:1251
#1 0x00000000009bd2c7 in rtl_check_failed_type2 (r=0x7ffff688cd40, n=<optimized out>, c1=<optimized out>, c2=<optimized out>, file=<optimized out>, line=<optimized out>, func=0x106ac48 <_ZZ7PATTERNP7rtx_defE12__FUNCTION__> "PATTERN") at [...]/gcc/rtl.c:802
#2 0x0000000000529ef3 in PATTERN (insn=<optimized out>) at [...]/gcc/rtl.h:1440
#3 0x00000000005e5a2b in PATTERN (insn=<optimized out>) at [...]/gcc/rtl.h:1440
#4 0x0000000000d08b96 in nvptx_propagate_unified (unified=0x7ffff688ccc0) at [...]/gcc/config/nvptx/nvptx.c:2299
#5 0x0000000000d093e7 in nvptx_split_blocks (map=map@entry=0x7fffffffcc40) at [...]/gcc/config/nvptx/nvptx.c:2428
#6 0x0000000000d0d08b in nvptx_reorg () at [...]/gcc/config/nvptx/nvptx.c:3840
#7 0x00000000009bb0ea in (anonymous namespace)::pass_machine_reorg::execute (this=<optimized out>) at [...]/gcc/reorg.c:3952
[...]
(gdb) frame 4
#4 0x0000000000d08b96 in nvptx_propagate_unified (unified=0x7ffff688ccc0) at [...]/gcc/config/nvptx/nvptx.c:2299
2299 pat = PATTERN (probe);
(gdb) print probe
$1 = (rtx_insn *) 0x7ffff688cd40
(gdb) call debug_rtx(probe)
(note 56 54 57 3 NOTE_INSN_DELETED)
>
> 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.
> --- 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);
These problems (both yours and mine) do not reproduce on trunk, right?
But I suppose these are still a latent, just waiting for a different test
case? Maybe this is a case to write an RTL-level test case? (Unless the
fix is deemed trivial enough to warrent spending time on this.)
Anyway, I don't know a lot about RTL, but the following patch does cure
this test case (now running other testing). Would you please check that,
and also whether nvptx_propagate_unified then still works as expected?
Is this patch OK (both for gomp-4_0-branch, and also for trunk?), or
should this rather use something like:
-if (!INSN_P (probe))
+if (NOTE_P (probe) && NOTE_KIND (probe) == NOTE_INSN_DELETED)
continue;
..., or something yet different?
--- gcc/config/nvptx/nvptx.c
+++ gcc/config/nvptx/nvptx.c
@@ -2286,7 +2286,7 @@ nvptx_propagate_unified (rtx_insn *unified)
{
rtx_insn *probe = unified;
rtx cond_reg = SET_DEST (PATTERN (unified));
- rtx pat;
+ rtx pat = NULL_RTX;
/* Find the comparison. (We could skip this and simply scan to he
blocks' terminating branch, if we didn't care for self
@@ -2294,14 +2294,17 @@ nvptx_propagate_unified (rtx_insn *unified)
for (;;)
{
probe = NEXT_INSN (probe);
+ if (!INSN_P (probe))
+ continue;
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) || !INSN_P (probe));
+ gcc_assert (NONJUMP_INSN_P (probe));
}
+ gcc_assert (pat != NULL_RTX);
rtx pred_reg = SET_DEST (pat);
/* Find the branch. */
Grüße
Thomas
^ permalink raw reply [flat|nested] 3+ messages in thread
* [gomp4, nvptx, committed] Fix --enable-checking=rtl assert in nvptx_propagate_unified
2017-07-06 7:16 ` Thomas Schwinge
@ 2017-07-16 18:59 ` Tom de Vries
0 siblings, 0 replies; 3+ messages in thread
From: Tom de Vries @ 2017-07-16 18:59 UTC (permalink / raw)
To: Thomas Schwinge, GCC Patches
[-- Attachment #1: Type: text/plain, Size: 5548 bytes --]
[ was: Re: [gomp4, nvptx, committed] Fix assert in nvptx_propagate_unified ]
On 07/06/2017 09:15 AM, Thomas Schwinge wrote:
> Hi Tom!
>
> On Fri, 30 Jun 2017 17:15:24 +0200, Tom de Vries <Tom_deVries@mentor.com> wrote:
>> with the openacc test-case in attached patch, I ran into an assert here:
>
> Using your test case, in my build with
> "--enable-checking=yes,df,fold,rtl", I already earlier run into an ICE...
>
Hi,
I see, thanks for letting me know.
Hmm, it looks like the build and test config with which I ran into this
problem only has the default checking (All my private build and test
configs use yes,rtl).
>> 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);
>
> ... here:
>
> [...]/libgomp/testsuite/libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c:19:9: internal compiler error: RTL check: expected elt 3 type 'e' or 'u', have '0' (rtx note) in PATTERN, at rtl.h:1440
>
> Breakpoint 2, internal_error (gmsgid=0x10cc840 "RTL check: expected elt %d type '%c' or '%c', have '%c' (rtx %s) in %s, at %s:%d") at [...]/gcc/diagnostic.c:1251
> 1251 {
> (gdb) bt
> #0 internal_error (gmsgid=0x10cc840 "RTL check: expected elt %d type '%c' or '%c', have '%c' (rtx %s) in %s, at %s:%d") at [...]/gcc/diagnostic.c:1251
> #1 0x00000000009bd2c7 in rtl_check_failed_type2 (r=0x7ffff688cd40, n=<optimized out>, c1=<optimized out>, c2=<optimized out>, file=<optimized out>, line=<optimized out>, func=0x106ac48 <_ZZ7PATTERNP7rtx_defE12__FUNCTION__> "PATTERN") at [...]/gcc/rtl.c:802
> #2 0x0000000000529ef3 in PATTERN (insn=<optimized out>) at [...]/gcc/rtl.h:1440
> #3 0x00000000005e5a2b in PATTERN (insn=<optimized out>) at [...]/gcc/rtl.h:1440
> #4 0x0000000000d08b96 in nvptx_propagate_unified (unified=0x7ffff688ccc0) at [...]/gcc/config/nvptx/nvptx.c:2299
> #5 0x0000000000d093e7 in nvptx_split_blocks (map=map@entry=0x7fffffffcc40) at [...]/gcc/config/nvptx/nvptx.c:2428
> #6 0x0000000000d0d08b in nvptx_reorg () at [...]/gcc/config/nvptx/nvptx.c:3840
> #7 0x00000000009bb0ea in (anonymous namespace)::pass_machine_reorg::execute (this=<optimized out>) at [...]/gcc/reorg.c:3952
> [...]
> (gdb) frame 4
> #4 0x0000000000d08b96 in nvptx_propagate_unified (unified=0x7ffff688ccc0) at [...]/gcc/config/nvptx/nvptx.c:2299
> 2299 pat = PATTERN (probe);
> (gdb) print probe
> $1 = (rtx_insn *) 0x7ffff688cd40
> (gdb) call debug_rtx(probe)
> (note 56 54 57 3 NOTE_INSN_DELETED)
>
>>
>> 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.
>
>> --- 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);
>
> These problems (both yours and mine) do not reproduce on trunk, right?
Correct, AFAICT nvptx_propagate_unified is something that was not ported
to trunk yet.
> But I suppose these are still a latent, just waiting for a different test
> case? Maybe this is a case to write an RTL-level test case? (Unless the
> fix is deemed trivial enough to warrent spending time on this.)
>
> Anyway, I don't know a lot about RTL, but the following patch does cure
> this test case (now running other testing). Would you please check that,
> and also whether nvptx_propagate_unified then still works as expected?
> Is this patch OK (both for gomp-4_0-branch, and also for trunk?), or
> should this rather use something like:
>
> -if (!INSN_P (probe))
> +if (NOTE_P (probe) && NOTE_KIND (probe) == NOTE_INSN_DELETED)
> continue;
>
> ..., or something yet different?
>
I took most of your patch, but used next_real_insn to skip over the note.
Tested on x86_64 with nvptx accelerator.
Committed to gomp-4_0-branch.
Thanks,
- Tom
[-- Attachment #2: 0001-Fix-enable-checking-rtl-assert-in-nvptx_propagate_unified.patch --]
[-- Type: text/x-patch, Size: 1330 bytes --]
Fix --enable-checking=rtl assert in nvptx_propagate_unified
2017-07-16 Tom de Vries <tom@codesourcery.com>
* config/nvptx/nvptx.c (nvptx_propagate_unified): Use next_real_insn to
find comparison insn. Assert that comparison is found.
---
gcc/config/nvptx/nvptx.c | 9 ++++++---
1 file changed, 6 insertions(+), 3 deletions(-)
diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index 69b5740..32420c3 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -2294,22 +2294,25 @@ nvptx_propagate_unified (rtx_insn *unified)
{
rtx_insn *probe = unified;
rtx cond_reg = SET_DEST (PATTERN (unified));
- rtx pat;
+ rtx pat = NULL_RTX;
/* 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);
+ probe = next_real_insn (probe);
+ if (!probe)
+ break;
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) || !INSN_P (probe));
+ gcc_assert (NONJUMP_INSN_P (probe));
}
+ gcc_assert (pat);
rtx pred_reg = SET_DEST (pat);
/* Find the branch. */
^ permalink raw reply [flat|nested] 3+ messages in thread
end of thread, other threads:[~2017-07-16 18:59 UTC | newest]
Thread overview: 3+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
[not found] <43470dd6-1f0b-7602-703a-60c8543c74e3@mentor.com>
2017-06-30 15:15 ` [gomp4, nvptx, committed] Fix assert in nvptx_propagate_unified Tom de Vries
2017-07-06 7:16 ` Thomas Schwinge
2017-07-16 18:59 ` [gomp4, nvptx, committed] Fix --enable-checking=rtl " 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).