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

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index c51aaf995c5e87694307d1a4a3c957e3fcc3b962..f547931b4e9d7b9a0506808f054dd25eff4eca11 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,8 @@
+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>
 

diff --git a/gcc/lto-streamer-in.c b/gcc/lto-streamer-in.c
index 299900ae3799797b202f66ad35b9d64b50a75836..2eb80511b919f82a77a392eb6258618a5dee6e06 100644 (file)
--- a/gcc/lto-streamer-in.c
+++ b/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;

diff --git a/gcc/lto-streamer-out.c b/gcc/lto-streamer-out.c
index 1b88115b2c98fe54282a221a706cb76f454eaf35..3ca88556638af7f442b7f8d650a0ba6eeab2a3aa 100644 (file)
--- a/gcc/lto-streamer-out.c
+++ b/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++)

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index ccdeaa71ffb9bda265461c6f215db5ddcd5b6297..3b60290b0bcb1a8e32edf406866c3ca5556660ab 100644 (file)
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,7 @@
+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.

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-type-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-type-1.c
new file mode 100644 (file)
index 0000000..5adfcec
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-type-1.c
@@ -0,0 +1,24 @@
+#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;
+}