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

gcc/ChangeLog

+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

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;

libgomp/ChangeLog

+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

libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c

+/* { 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;
+}