| Message ID | 87k2t8le94.fsf@kepler.schwinge.homeip.net |
|---|---|
| State | New |
| Headers | show |
Ok. Richard. On 8/6/15, Thomas Schwinge <thomas@codesourcery.com> wrote: > Hi! > > On Wed, 5 Aug 2015 15:10:40 +0100, "David Sherwood" <david.sherwood@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:thomas@codesourcery.com] >> > On Wed, 5 Aug 2015 11:18:32 +0100, David Sherwood >> > <david.sherwood@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 <iverbin@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 >> > > > > > <iverbin@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 <thomas@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 >
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; +}