From 96a2d174ccfd5b1cdb746e79468dd926ad26718f Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Mon, 10 Aug 2015 17:22:30 +0200 Subject: [PATCH] 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 | 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(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/vector-type-1.c diff --git a/gcc/ChangeLog b/gcc/ChangeLog index c51aaf995c5..f547931b4e9 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,8 @@ +2015-08-10 Thomas Schwinge + + * lto-streamer-in.c (lto_input_mode_table): Adjust to + GET_MODE_INNER changes. + 2015-08-10 Thomas Schwinge Ilya Verbin diff --git a/gcc/lto-streamer-in.c b/gcc/lto-streamer-in.c index 299900ae379..2eb80511b91 100644 --- 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 1b88115b2c9..3ca88556638 100644 --- 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 ccdeaa71ffb..3b60290b0bc 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,7 @@ +2015-08-10 Thomas Schwinge + + * testsuite/libgomp.oacc-c-c++-common/vector-type-1.c: New file. + 2015-08-03 Nathan Sidwell * 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 index 00000000000..5adfcecd641 --- /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; +} -- 2.30.2