nvptx.c (nvptx_declare_function_name): Round frame size to DImode boundary.

	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.

From-SVN: r239086

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 3da9eb70fff2a026f430e2037e65f90b7e92514f..2223290bd53469d4a56cf024be0aab2adc1fdf76 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,9 @@
+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.
+
 2016-08-03  Alan Modra  <amodra@gmail.com>
 
        * config/rs6000/rs6000.h (SLOW_UNALIGNED_ACCESS): Make scalar

diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index 6c78699d75f6bdae0dc25bd6452352a8eb065bfc..4d87ead20f39aeec039d8f94b73104856f3ff4c6 100644 (file)
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -1037,11 +1037,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 ();
@@ -3266,8 +3269,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 a/libgomp/ChangeLog b/libgomp/ChangeLog
index ebb4bc93369552cf816545eaa9516d3e73e11baf..850188fb660c8c640e8f8ce0a06e6ad416b5048a 100644 (file)
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,7 @@
+2016-08-03  Nathan Sidwell  <nathan@codesourcery.com>
+
+       * testsuite/libgomp.oacc-c-c++-common/crash-1.c: New.
+
 2016-07-15  Cesar Philippidis  <cesar@codesourcery.com>
 
        * testsuite/libgomp.oacc-c-c++-common/zero_length_subarrays.c: New

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c
new file mode 100644 (file)
index 0000000..a75a817
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c
@@ -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;
+}