Hi!
On Wed, 3 Aug 2016 13:30:10 -0400, Nathan Sidwell <nat...@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 <nat...@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 <nat...@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 <tho...@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 <tho...@codesourcery.com>
+
+ Backport trunk r239086:
+ 2016-08-03 Nathan Sidwell <nat...@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 <jasonw...@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 <tho...@codesourcery.com>
+
+ Backport trunk r239125:
+ 2016-08-04 Thomas Schwinge <tho...@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 <nat...@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
