Hi! On Sat, 8 Aug 2015 07:25:42 +0200, Richard Biener <richard.guent...@gmail.com> wrote: > Ok.
Committed in r226758 and r226759:
commit 7231f6b984806cceb30cacf0e79f8f5ae7a68803
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Mon Aug 10 15:22:24 2015 +0000
Correctly advance iterator in offloading machine mode stream reading
gcc/
* lto-streamer-in.c (lto_input_mode_table): Correctly advance
iterator.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@226758
138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog | 6 ++++++
gcc/lto-streamer-in.c | 2 +-
2 files changed, 7 insertions(+), 1 deletion(-)
diff --git gcc/ChangeLog gcc/ChangeLog
index f103d41..c51aaf9 100644
--- gcc/ChangeLog
+++ gcc/ChangeLog
@@ -1,3 +1,9 @@
+2015-08-10 Thomas Schwinge <tho...@codesourcery.com>
+ Ilya Verbin <ilya.ver...@intel.com>
+
+ * lto-streamer-in.c (lto_input_mode_table): Correctly advance
+ iterator.
+
2015-08-09 Manuel López-Ibáñez <m...@gcc.gnu.org>
* doc/options.texi (EnabledBy): Document that the argument must be
diff --git gcc/lto-streamer-in.c gcc/lto-streamer-in.c
index a56d3f3..299900a 100644
--- gcc/lto-streamer-in.c
+++ gcc/lto-streamer-in.c
@@ -1573,7 +1573,7 @@ lto_input_mode_table (struct lto_file_decl_data
*file_data)
for (machine_mode mr = pass ? VOIDmode
: GET_CLASS_NARROWEST_MODE (mclass);
pass ? mr < MAX_MACHINE_MODE : mr != VOIDmode;
- pass ? mr = (machine_mode) (m + 1)
+ pass ? mr = (machine_mode) (mr + 1)
: mr = GET_MODE_WIDER_MODE (mr))
if (GET_MODE_CLASS (mr) != mclass
|| GET_MODE_SIZE (mr) != size
commit b308f4a0d03e67bdaf3f43416cfbd360db957a29
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Mon Aug 10 15:22:30 2015 +0000
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.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@226759
138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog | 5 ++++
gcc/lto-streamer-in.c | 8 ++++---
gcc/lto-streamer-out.c | 4 ++--
libgomp/ChangeLog | 4 ++++
.../libgomp.oacc-c-c++-common/vector-type-1.c | 24 ++++++++++++++++++++
5 files changed, 40 insertions(+), 5 deletions(-)
diff --git gcc/ChangeLog gcc/ChangeLog
index c51aaf9..f547931 100644
--- gcc/ChangeLog
+++ gcc/ChangeLog
@@ -1,4 +1,9 @@
2015-08-10 Thomas Schwinge <tho...@codesourcery.com>
+
+ * lto-streamer-in.c (lto_input_mode_table): Adjust to
+ GET_MODE_INNER changes.
+
+2015-08-10 Thomas Schwinge <tho...@codesourcery.com>
Ilya Verbin <ilya.ver...@intel.com>
* lto-streamer-in.c (lto_input_mode_table): Correctly advance
diff --git gcc/lto-streamer-in.c gcc/lto-streamer-in.c
index 299900a..2eb8051 100644
--- gcc/lto-streamer-in.c
+++ 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 gcc/lto-streamer-out.c gcc/lto-streamer-out.c
index 1b88115..3ca8855 100644
--- gcc/lto-streamer-out.c
+++ 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 libgomp/ChangeLog libgomp/ChangeLog
index ccdeaa7..3b60290 100644
--- libgomp/ChangeLog
+++ libgomp/ChangeLog
@@ -1,3 +1,7 @@
+2015-08-10 Thomas Schwinge <tho...@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c-c++-common/vector-type-1.c: New file.
+
2015-08-03 Nathan Sidwell <nat...@codesourcery.com>
* plugin/plugin-nvptx.c: Don't include dlfcn.h.
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/vector-type-1.c
libgomp/testsuite/libgomp.oacc-c-c++-common/vector-type-1.c
new file mode 100644
index 0000000..5adfcec
--- /dev/null
+++ 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;
+}
Grüße,
Thomas
pgplDDeUTCVhD.pgp
Description: PGP signature
