From: Nathan Sidwell <nathan@codesourcery.com>
Date: Wed, 3 Aug 2016 17:26:51 +0000 (+0000)
Subject: nvptx.c (nvptx_declare_function_name): Round frame size to DImode boundary.
X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=c1311c86c0c1057b43146c8817dfe453c7a88fd4;p=gcc.git

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 3da9eb70fff..2223290bd53 100644
--- 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 6c78699d75f..4d87ead20f3 100644
--- 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 ebb4bc93369..850188fb660 100644
--- 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
index 00000000000..a75a81765c6
--- /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;
+}