* [PTX] fix worker propagation ICE
@ 2016-08-03 17:30 Nathan Sidwell
2016-08-04 13:52 ` Thomas Schwinge
2017-01-10 12:09 ` Thomas Schwinge
0 siblings, 2 replies; 3+ messages in thread
From: Nathan Sidwell @ 2016-08-03 17:30 UTC (permalink / raw)
To: GCC Patches
[-- Attachment #1: Type: text/plain, Size: 597 bytes --]
The PTX backend could ice when generating a state propagation sequence entering
partitioned execution. Although the stack frame is DImode aligned, nothing
actually rounds the size up consistent with that. That meant we could encounter
frames that were not a DImode multiple in size. Which broke the assert checking
that.
Rather than faff around trying to copy just the extra bit on the end of such a
frame, I changed the frame emission to round the size up, and adjust the
propagation machinery likewise. (Mostly one gets frames when not optimizing
anyway).
Applied to trunk & gomp4.
[-- Attachment #2: ptx-ice.patch --]
[-- Type: text/x-patch, Size: 2325 bytes --]
2016-08-03 Nathan Sidwell <nathan@codesourcery.com>
gcc/
* config/nvptx/nvptx.c (nvptx_declare_function_name): Round frame
size to DImode boundary.
(nvptx_propagate): Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/crash-1.c: New.
Index: gcc/config/nvptx/nvptx.c
===================================================================
--- gcc/config/nvptx/nvptx.c (revision 239084)
+++ gcc/config/nvptx/nvptx.c (working copy)
@@ -999,11 +999,14 @@ nvptx_declare_function_name (FILE *file,
init_frame (file, STACK_POINTER_REGNUM,
UNITS_PER_WORD, crtl->outgoing_args_size);
- /* Declare a local variable for the frame. */
+ /* Declare a local variable for the frame. Force its size to be
+ DImode-compatible. */
HOST_WIDE_INT sz = get_frame_size ();
if (sz || cfun->machine->has_chain)
init_frame (file, FRAME_POINTER_REGNUM,
- crtl->stack_alignment_needed / BITS_PER_UNIT, sz);
+ crtl->stack_alignment_needed / BITS_PER_UNIT,
+ (sz + GET_MODE_SIZE (DImode) - 1)
+ & ~(HOST_WIDE_INT)(GET_MODE_SIZE (DImode) - 1));
/* Declare the pseudos we have as ptx registers. */
int maxregs = max_reg_num ();
@@ -3222,8 +3225,9 @@ nvptx_propagate (basic_block block, rtx_
rtx pred = NULL_RTX;
rtx_code_label *label = NULL;
- gcc_assert (!(fs & (GET_MODE_SIZE (DImode) - 1)));
- fs /= GET_MODE_SIZE (DImode);
+ /* The frame size might not be DImode compatible, but the frame
+ array's declaration will be. So it's ok to round up here. */
+ fs = (fs + GET_MODE_SIZE (DImode) - 1) / GET_MODE_SIZE (DImode);
/* Detect single iteration loop. */
if (fs == 1)
fs = 0;
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c (nonexistent)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c (working copy)
@@ -0,0 +1,28 @@
+/* { dg-do compile } */
+/* { dg-options "-O0" } */
+
+/* ICEd in nvptx backend due to unexpected frame size. */
+#pragma acc routine worker
+void
+worker_matmul (int *c, int i)
+{
+ int j;
+
+#pragma acc loop
+ for (j = 0; j < 4; j++)
+ c[j] = j;
+}
+
+
+int
+main ()
+{
+ int c[4];
+
+#pragma acc parallel
+ {
+ worker_matmul (c, 0);
+ }
+
+ return 0;
+}
^ permalink raw reply [flat|nested] 3+ messages in thread
* Re: [PTX] fix worker propagation ICE
2016-08-03 17:30 [PTX] fix worker propagation ICE Nathan Sidwell
@ 2016-08-04 13:52 ` Thomas Schwinge
2017-01-10 12:09 ` Thomas Schwinge
1 sibling, 0 replies; 3+ messages in thread
From: Thomas Schwinge @ 2016-08-04 13:52 UTC (permalink / raw)
To: Nathan Sidwell, GCC Patches
[-- Attachment #1: Type: text/plain, Size: 4029 bytes --]
Hi!
On Wed, 3 Aug 2016 13:30:10 -0400, Nathan Sidwell <nathan@acm.org> wrote:
> --- libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c (nonexistent)
> +++ libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c (working copy)
> @@ -0,0 +1,28 @@
> +/* { dg-do compile } */
> +/* { dg-options "-O0" } */
Offloading compilation happens at link time not compile time, and in
OpenACC libgomp testing, we're doing a limited set of torture testing
(-O0, -O2), so no point in hardcoding -O0 here.
As obvious, committed to trunk in r239125:
commit ec8d61ace153843dcaaba86ad926f384a4affee3
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Thu Aug 4 13:34:57 2016 +0000
Make libgomp.oacc-c-c++-common/crash-1.c a "link" test, and don't hardcode -O0
libgomp/
* testsuite/libgomp.oacc-c-c++-common/crash-1.c: Make it a "link"
test, and don't hardcode -O0.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@239125 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog | 5 +++++
libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c | 5 ++---
2 files changed, 7 insertions(+), 3 deletions(-)
diff --git libgomp/ChangeLog libgomp/ChangeLog
index 850188f..cc76b7b 100644
--- libgomp/ChangeLog
+++ libgomp/ChangeLog
@@ -1,3 +1,8 @@
+2016-08-04 Thomas Schwinge <thomas@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c-c++-common/crash-1.c: Make it a "link"
+ test, and don't hardcode -O0.
+
2016-08-03 Nathan Sidwell <nathan@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/crash-1.c: New.
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c
index a75a817..dcf1485 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c
@@ -1,7 +1,6 @@
-/* { dg-do compile } */
-/* { dg-options "-O0" } */
+/* { dg-do link } */
-/* ICEd in nvptx backend due to unexpected frame size. */
+/* For -O0, ICEd in nvptx backend due to unexpected frame size. */
#pragma acc routine worker
void
worker_matmul (int *c, int i)
Backported to gomp-4_0-branch in r239129:
commit bab445509b917c582a53834599f614ce2c29ff36
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Thu Aug 4 13:49:15 2016 +0000
Make libgomp.oacc-c-c++-common/crash-1.c a "link" test, and don't hardcode -O0
Backport trunk r239125:
libgomp/
* testsuite/libgomp.oacc-c-c++-common/crash-1.c: Make it a "link"
test, and don't hardcode -O0.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@239129 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog.gomp | 6 ++++++
libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c | 5 ++---
2 files changed, 8 insertions(+), 3 deletions(-)
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 7a7d859..4320237 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,3 +1,9 @@
+2016-08-04 Thomas Schwinge <thomas@codesourcery.com>
+
+ Backport trunk r239125:
+ * testsuite/libgomp.oacc-c-c++-common/crash-1.c: Make it a "link"
+ test, and don't hardcode -O0.
+
2016-08-03 Nathan Sidwell <nathan@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/crash-1.c: New.
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c
index a75a817..dcf1485 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c
@@ -1,7 +1,6 @@
-/* { dg-do compile } */
-/* { dg-options "-O0" } */
+/* { dg-do link } */
-/* ICEd in nvptx backend due to unexpected frame size. */
+/* For -O0, ICEd in nvptx backend due to unexpected frame size. */
#pragma acc routine worker
void
worker_matmul (int *c, int i)
Grüße
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 3+ messages in thread
* Re: [PTX] fix worker propagation ICE
2016-08-03 17:30 [PTX] fix worker propagation ICE Nathan Sidwell
2016-08-04 13:52 ` Thomas Schwinge
@ 2017-01-10 12:09 ` Thomas Schwinge
1 sibling, 0 replies; 3+ messages in thread
From: Thomas Schwinge @ 2017-01-10 12:09 UTC (permalink / raw)
To: GCC Patches
Hi!
On Wed, 3 Aug 2016 13:30:10 -0400, Nathan Sidwell <nathan@acm.org> wrote:
> The PTX backend could ice when generating a state propagation sequence entering
> partitioned execution. [...]
To fix the same problem, applying the same changes, committed the
following to gcc-6-branch in r244261:
commit 71f5a3c5007ba5bd398d978a56640b7163530ab5
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Tue Jan 10 12:07:13 2017 +0000
[PTX] fix worker propagation ICE
Backport trunk r239086:
gcc/
2016-08-03 Nathan Sidwell <nathan@codesourcery.com>
* config/nvptx/nvptx.c (nvptx_declare_function_name): Round frame
size to DImode boundary.
(nvptx_propagate): Likewise.
libgomp/
2016-08-03 Nathan Sidwell <nathan@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/crash-1.c: New.
Backport trunk r239125 'Make libgomp.oacc-c-c++-common/crash-1.c a "link" test,
and don't hardcode -O0':
libgomp/
2016-08-04 Thomas Schwinge <thomas@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/crash-1.c: Make it a "link"
test, and don't hardcode -O0.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gcc-6-branch@244261 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog | 9 ++++++++
gcc/config/nvptx/nvptx.c | 12 ++++++----
libgomp/ChangeLog | 13 +++++++++++
.../testsuite/libgomp.oacc-c-c++-common/crash-1.c | 27 ++++++++++++++++++++++
4 files changed, 57 insertions(+), 4 deletions(-)
diff --git gcc/ChangeLog gcc/ChangeLog
index 71b0742..86b7f1b 100644
--- gcc/ChangeLog
+++ gcc/ChangeLog
@@ -1,3 +1,12 @@
+2017-01-10 Thomas Schwinge <thomas@codesourcery.com>
+
+ Backport trunk r239086:
+ 2016-08-03 Nathan Sidwell <nathan@codesourcery.com>
+
+ * config/nvptx/nvptx.c (nvptx_declare_function_name): Round frame
+ size to DImode boundary.
+ (nvptx_propagate): Likewise.
+
2017-01-10 Chung-Ju Wu <jasonwucj@gmail.com>
Backport from mainline
diff --git gcc/config/nvptx/nvptx.c gcc/config/nvptx/nvptx.c
index a6c90b6..2262005 100644
--- gcc/config/nvptx/nvptx.c
+++ gcc/config/nvptx/nvptx.c
@@ -989,11 +989,14 @@ nvptx_declare_function_name (FILE *file, const char *name, const_tree decl)
init_frame (file, STACK_POINTER_REGNUM,
UNITS_PER_WORD, crtl->outgoing_args_size);
- /* Declare a local variable for the frame. */
+ /* Declare a local variable for the frame. Force its size to be
+ DImode-compatible. */
HOST_WIDE_INT sz = get_frame_size ();
if (sz || cfun->machine->has_chain)
init_frame (file, FRAME_POINTER_REGNUM,
- crtl->stack_alignment_needed / BITS_PER_UNIT, sz);
+ crtl->stack_alignment_needed / BITS_PER_UNIT,
+ (sz + GET_MODE_SIZE (DImode) - 1)
+ & ~(HOST_WIDE_INT)(GET_MODE_SIZE (DImode) - 1));
/* Declare the pseudos we have as ptx registers. */
int maxregs = max_reg_num ();
@@ -3212,8 +3215,9 @@ nvptx_propagate (basic_block block, rtx_insn *insn, propagate_mask rw,
rtx pred = NULL_RTX;
rtx_code_label *label = NULL;
- gcc_assert (!(fs & (GET_MODE_SIZE (DImode) - 1)));
- fs /= GET_MODE_SIZE (DImode);
+ /* The frame size might not be DImode compatible, but the frame
+ array's declaration will be. So it's ok to round up here. */
+ fs = (fs + GET_MODE_SIZE (DImode) - 1) / GET_MODE_SIZE (DImode);
/* Detect single iteration loop. */
if (fs == 1)
fs = 0;
diff --git libgomp/ChangeLog libgomp/ChangeLog
index 8841636..2e5f73b 100644
--- libgomp/ChangeLog
+++ libgomp/ChangeLog
@@ -1,3 +1,16 @@
+2017-01-10 Thomas Schwinge <thomas@codesourcery.com>
+
+ Backport trunk r239125:
+ 2016-08-04 Thomas Schwinge <thomas@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c-c++-common/crash-1.c: Make it a "link"
+ test, and don't hardcode -O0.
+
+ Backport trunk r239086:
+ 2016-08-03 Nathan Sidwell <nathan@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c-c++-common/crash-1.c: New.
+
2016-12-21 Release Manager
* GCC 6.3.0 released.
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c
new file mode 100644
index 0000000..dcf1485
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c
@@ -0,0 +1,27 @@
+/* { dg-do link } */
+
+/* For -O0, ICEd in nvptx backend due to unexpected frame size. */
+#pragma acc routine worker
+void
+worker_matmul (int *c, int i)
+{
+ int j;
+
+#pragma acc loop
+ for (j = 0; j < 4; j++)
+ c[j] = j;
+}
+
+
+int
+main ()
+{
+ int c[4];
+
+#pragma acc parallel
+ {
+ worker_matmul (c, 0);
+ }
+
+ return 0;
+}
Grüße
Thomas
^ permalink raw reply [flat|nested] 3+ messages in thread
end of thread, other threads:[~2017-01-10 12:09 UTC | newest]
Thread overview: 3+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2016-08-03 17:30 [PTX] fix worker propagation ICE Nathan Sidwell
2016-08-04 13:52 ` Thomas Schwinge
2017-01-10 12:09 ` Thomas Schwinge
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).