Fix offloading machine mode stream reading

... in context of the GET_MODE_INNER changes applied in r226328.

	gcc/
	* lto-streamer-in.c (lto_input_mode_table): Adjust to
	GET_MODE_INNER changes.
	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/vector-type-1.c: New file.

From-SVN: r226759

gcc/ChangeLog

 2015-08-10  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* lto-streamer-in.c (lto_input_mode_table): Adjust to
+	GET_MODE_INNER changes.
+
+2015-08-10  Thomas Schwinge  <thomas@codesourcery.com>
 	    Ilya Verbin  <ilya.verbin@intel.com>
 
 	* lto-streamer-in.c (lto_input_mode_table): Correctly advance

gcc/lto-streamer-in.c

@@ -1544,7 +1544,7 @@ lto_input_mode_table (struct lto_file_decl_data *file_data)
 	= bp_unpack_enum (&bp, mode_class, MAX_MODE_CLASS);
       unsigned int size = bp_unpack_value (&bp, 8);
       unsigned int prec = bp_unpack_value (&bp, 16);
-      machine_mode inner = (machine_mode) table[bp_unpack_value (&bp, 8)];
+      machine_mode inner = (machine_mode) bp_unpack_value (&bp, 8);
       unsigned int nunits = bp_unpack_value (&bp, 8);
       unsigned int ibit = 0, fbit = 0;
       unsigned int real_fmt_len = 0;
@@ -1578,7 +1578,9 @@ lto_input_mode_table (struct lto_file_decl_data *file_data)
 	  if (GET_MODE_CLASS (mr) != mclass
 	      || GET_MODE_SIZE (mr) != size
 	      || GET_MODE_PRECISION (mr) != prec
-	      || GET_MODE_INNER (mr) != inner
+	      || (inner == m
+		  ? GET_MODE_INNER (mr) != mr
+		  : GET_MODE_INNER (mr) != table[(int) inner])
 	      || GET_MODE_IBIT (mr) != ibit
 	      || GET_MODE_FBIT (mr) != fbit
 	      || GET_MODE_NUNITS (mr) != nunits)
@@ -1606,7 +1608,7 @@ lto_input_mode_table (struct lto_file_decl_data *file_data)
 	    case MODE_VECTOR_UACCUM:
 	      /* For unsupported vector modes just use BLKmode,
 		 if the scalar mode is supported.  */
-	      if (inner != VOIDmode)
+	      if (table[(int) inner] != VOIDmode)
 		{
 		  table[m] = BLKmode;
 		  break;

gcc/lto-streamer-out.c

@@ -2676,7 +2676,7 @@ lto_write_mode_table (void)
   ob = create_output_block (LTO_section_mode_table);
   bitpack_d bp = bitpack_create (ob->main_stream);
 
-  /* Ensure that for GET_MODE_INNER (m) != VOIDmode we have
+  /* Ensure that for GET_MODE_INNER (m) != m we have
      also the inner mode marked.  */
   for (int i = 0; i < (int) MAX_MACHINE_MODE; i++)
     if (streamer_mode_table[i])
@@ -2685,7 +2685,7 @@ lto_write_mode_table (void)
 	if (GET_MODE_INNER (m) != m)
 	  streamer_mode_table[(int) GET_MODE_INNER (m)] = 1;
       }
-  /* First stream modes that have GET_MODE_INNER (m) == VOIDmode,
+  /* First stream modes that have GET_MODE_INNER (m) == m,
      so that we can refer to them afterwards.  */
   for (int pass = 0; pass < 2; pass++)
     for (int i = 0; i < (int) MAX_MACHINE_MODE; i++)

libgomp/ChangeLog

+2015-08-10  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* testsuite/libgomp.oacc-c-c++-common/vector-type-1.c: New file.
+
 2015-08-03  Nathan Sidwell  <nathan@codesourcery.com>
 
 	* plugin/plugin-nvptx.c: Don't include dlfcn.h.

libgomp/testsuite/libgomp.oacc-c-c++-common/vector-type-1.c

+#define vector __attribute__ ((vector_size (4 * sizeof(int))))
+
+int main(void)
+{
+  vector int vi = { 12, -34, -56, 78 };
+
+#pragma acc parallel copy(vi)
+  {
+    if (vi[0] != 12
+	|| vi[1] != -34
+	|| vi[2] != -56
+	|| vi[3] != 78)
+      __builtin_abort();
+    vector int vi_ = { -21, -43, 65, 87 };
+    vi = vi_;
+  }
+  if (vi[0] != -21
+      || vi[1] != -43
+      || vi[2] != 65
+      || vi[3] != 87)
+    __builtin_abort();
+
+  return 0;
+}