Hi! On Wed, 5 Aug 2015 15:10:40 +0100, "David Sherwood" <david.sherw...@arm.com> wrote: > In lto_input_mode_table there is the following line of code: > > machine_mode inner = (machine_mode) table[bp_unpack_value (&bp, 8)]; > > Is this right? In lto_write_mode_table this inner mode is written out > explicitly > into the stream already, so do we just need this instead? > > machine_mode inner = (machine_mode) bp_unpack_value (&bp, 8); > > It's possible I'm misunderstanding the code somehow though ...
The idea here is to translate between the machine mode IDs used by the target and the offloading compiler(s), whence the table lookup, or table construction in the first place. But as I said yesterday, you gave me a clue where to look, and the problem is that given your GET_MODE_INNER change: > > From: Thomas Schwinge [mailto:tho...@codesourcery.com] > > On Wed, 5 Aug 2015 11:18:32 +0100, David Sherwood <david.sherw...@arm.com> > > wrote: > > > I recently changed GET_MODE_INNER (m) > > > to return 'm' itself if there is no inner mode ... the following code from gcc/lto-streamer-in.c:lto_input_mode_table: > > > > On Tue, 4 Aug 2015 16:06:23 +0300, Ilya Verbin <iver...@gmail.com> > > > > wrote: > > > > > On Tue, Aug 04, 2015 at 14:35:11 +0200, Thomas Schwinge wrote: > > > > > > On Fri, 31 Jul 2015 20:13:02 +0300, Ilya Verbin <iver...@gmail.com> > > > > > > wrote: > > > > > > > On Fri, Jul 31, 2015 at 18:59:59 +0200, Jakub Jelinek wrote: > > > > > > > > > > On Wed, Feb 18, 2015 at 11:00:35 +0100, Jakub Jelinek wrote: > > > > > > > > > > + /* First search just the GET_CLASS_NARROWEST_MODE to > > > > > > > > > > wider modes, > > > > > > > > > > + if not found, fallback to all modes. */ > > > > > > > > > > + int pass; > > > > > > > > > > + for (pass = 0; pass < 2; pass++) > > > > > > > > > > + for (machine_mode mr = pass ? VOIDmode > > > > > > > > > > + : GET_CLASS_NARROWEST_MODE > > > > > > > > > > (mclass); > > > > > > > > > > + pass ? mr < MAX_MACHINE_MODE : mr != VOIDmode; > > > > > > > > > > + pass ? mr = (machine_mode) (m + 1) > > > > > > > > > > + : mr = GET_MODE_WIDER_MODE (mr)) > > > > > > > > > > + if (GET_MODE_CLASS (mr) != mclass > > > > > > > > > > + || GET_MODE_SIZE (mr) != size > > > > > > > > > > + || GET_MODE_PRECISION (mr) != prec > > > > > > > > > > + || GET_MODE_INNER (mr) != inner > > > > > > > > > > + || GET_MODE_IBIT (mr) != ibit > > > > > > > > > > + || GET_MODE_FBIT (mr) != fbit > > > > > > > > > > + || GET_MODE_NUNITS (mr) != nunits) > > > > > > > > > > + continue; ... no longer does the right thing in the »GET_MODE_INNER (mr) != inner« comparison. > > > > I'm trying, but I cannot claim yet to really understand this > > > > mode streaming code... ;-) Now that I do... > > > > But, with the producer > > > > gcc/lto-streamer-out.c:lto_write_mode_table having been changed, does > > > > maybe the consumer gcc/lto-streamer-in.c:lto_input_mode_table also need > > > > to be updated accordingly? > > > > > > > > For reference, David's change to > > > > gcc/lto-streamer-out.c:lto_write_mode_table: > > > > > > > > @@ -2679,23 +2679,23 @@ lto_write_mode_table (void) > > > > /* Ensure that for GET_MODE_INNER (m) != VOIDmode we have > > > > also the inner mode marked. */ > > > > for (int i = 0; i < (int) MAX_MACHINE_MODE; i++) > > > > if (streamer_mode_table[i]) > > > > { > > > > machine_mode m = (machine_mode) i; > > > > - if (GET_MODE_INNER (m) != VOIDmode) > > > > + 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, > > > > 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++) > > > > if (streamer_mode_table[i] && i != (int) VOIDmode && i != (int) > > > > BLKmode) > > > > { > > > > machine_mode m = (machine_mode) i; > > > > - if ((GET_MODE_INNER (m) == VOIDmode) ^ (pass == 0)) > > > > + if ((GET_MODE_INNER (m) == m) ^ (pass == 0)) > > > > continue; > > > > bp_pack_value (&bp, m, 8); > > > > bp_pack_enum (&bp, mode_class, MAX_MODE_CLASS, GET_MODE_CLASS > > > > (m)); > > > > bp_pack_value (&bp, GET_MODE_SIZE (m), 8); > > > > bp_pack_value (&bp, GET_MODE_PRECISION (m), 16); > > > > bp_pack_value (&bp, GET_MODE_INNER (m), 8); ... I came up with the following patch to fix the offloading machine mode stream reading. OK to commit? commit 45264b009e988298fddab5417e12d36e2edeeb49 Author: Thomas Schwinge <tho...@codesourcery.com> Date: Thu Aug 6 12:00:01 2015 +0200 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. --- gcc/lto-streamer-in.c | 8 ++++--- gcc/lto-streamer-out.c | 4 ++-- .../libgomp.oacc-c-c++-common/vector-type-1.c | 24 ++++++++++++++++++++ 3 files changed, 31 insertions(+), 5 deletions(-) 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; > > > > (Also, the source code comments need to be updated?) 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/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
signature.asc
Description: PGP signature
