| Message ID | ea04f184-d7cb-32b6-032d-8ede3d604d0a@codesourcery.com |
|---|---|
| State | New |
| Headers | show |
| Series | [LTO] PR 86416 – improve lto1 diagnostic if a mode does not exist (esp. for offloading targets) | expand |
*ping* On 12/13/19 3:34 PM, Tobias Burnus wrote: > As long as one compiles for a single target, the message is unlikely > to appear. > > However, when compiling for offloading, the modes supported on the target > 'host' and on the target 'device' can be different. In particular, > 'long double' (when larger than double) and '__float128' might not be > available on the device. > > This gives currently errors like the following (see PR, comment 0): > >> lto1: fatal error: unsupported mode TF >> compilation terminated. >> mkoffload: fatal error: >> x86_64-pc-linux-gnu-accel-nvptx-none-gcc returned 1 exit status > > While the device target is hidden in > 'x86_64-pc-linux-gnu-accel-nvptx-none-gcc', it might make more sense > to add it more prominently. Additionally, an average user does > not understand what 'TF' or 'XF' means. > > Solutions: > (A) Add the target to the output > (B) Add a better description for the error > > I did both in the attached patch, giving: > lto1: fatal error: nvptx-none - 80-bit floating-point numbers > unsupported (mode 'XF') > lto1: fatal error: nvptx-none - 128-bit floating-point numbers > unsupported (mode 'TF') > > * (A) should be fine, I think. > > * But I am not 100% happy with (B). I think as interim solution, > it is acceptable as XF/TF are well defined and probably the most > common problem. — Alternatively, one could only commit (A) or > solve it more properly (how?). > > * If, e.g., 'long long' or 'integer(kind=16)' are used, only a > generic message is printed. A message such as "'__float128' not > supported" or "'real(kind=10)' not supported" is more helpful and > supporting all modes and not cherry picking those two would be > useful as well. > > The question is how to pass this information to lto-streamer-in.c; > it is available as TYPE_NAME – and, with debugging turned on, this > also gets passed on, but is also not be readily available in > lto_input_mode_table. – Suggestions? > > > Build on x86-64-gnu-linux and tested without offloading and with nvptx > offloading. > > Tobias >
On Fri, Dec 13, 2019 at 03:34:56PM +0100, Tobias Burnus wrote: > --- a/gcc/lto-streamer-in.c > +++ b/gcc/lto-streamer-in.c > @@ -1700,7 +1700,19 @@ lto_input_mode_table (struct lto_file_decl_data *file_data) > } > /* FALLTHRU */ > default: > - fatal_error (UNKNOWN_LOCATION, "unsupported mode %qs", mname); > + /* For offloading-target compilions, this is a user-facing > + message. See also target.def and machmode.def. */ > + if (strcmp (mname, "XF") == 0) > + fatal_error (UNKNOWN_LOCATION, > + "%s - 80-bit floating-point numbers unsupported " > + "(mode %qs)", TARGET_MACHINE, mname); > + else if (strcmp (mname, "TF") == 0) > + fatal_error (UNKNOWN_LOCATION, > + "%s - 128-bit floating-point numbers unsupported " > + "(mode %qs)", TARGET_MACHINE, mname); I don't like the above, it looks like a gross hack for two modes that are often used on x86_64, but as soon as you e.g. use powerpc64 KFmode instead or many others, you'll be in trouble. I'd say let lto_output_mode_table stream some string with a description of the mode (perhaps for offloading streaming only), whether that would be by looking e.g. through standard lang specific types, finding if they have such mode and printing the name of the type, or trying to describe the mode in words somehow, and then print that string during lto_input_mode_table. Jakub
Hi Jakub, thanks for the pointers; I was also not happy about part (B). The mode gets written at lto-streamer-out.c's lto_write_mode_table, which is already called with 'if (lto_stream_offload_p)'. — Looking at places where the name gets constructed, I realized that required information is already written to the mode table: The precision (for float, to distinguish between 80bit and 128bit floats) and the size. One now has a readable error message for float/complex/decimal-float and int(eger) modes – and the generic one for the rest; as vector get converted to their base type, if not found, this give a readable message for all common cases. Hence, one can write a much cleaner patch which is about as long as the hack but not hackish :-) Bootstrapped + regtested on x86-64-gnu-linux w/o offloading and with nvptx offloading. OK? Thanks, Tobias
On Wed, Dec 18, 2019 at 05:39:51PM +0100, Tobias Burnus wrote: > Hence, one can write a much cleaner patch which is about as long as the hack > but not hackish :-) > > Bootstrapped + regtested on x86-64-gnu-linux w/o offloading and with nvptx > offloading. > OK? LGTM. > 2019-12-18 Tobias Burnus <tobias@codesourcery.com> > > PR middle-end/86416 > * Makefile.in (CFLAGS-lto-streamer-in.o): Pass target_noncanonical on. > * lto-streamer-in.c (lto_input_mode_table): Improve unsupported-mode > diagnostic. > > PR middle-end/86416 > * testsuite/libgomp.c/pr86416-1.c: New. > * testsuite/libgomp.c/pr86416-2.c: New. Jakub
PR middle-end/86416 * Makefile.in (CFLAGS-lto-streamer-in.o): Pass target_noncanonical on. * lto-streamer-in.c (lto_input_mode_table): Use it; add special error diagnostic for missing XF and TF modes. PR middle-end/86416 * testsuite/libgomp.c/pr86416-1.c: New. * testsuite/libgomp.c/pr86416-2.c: New. diff --git a/gcc/Makefile.in b/gcc/Makefile.in index 6b857bd75de..657488d416b 100644 --- a/gcc/Makefile.in +++ b/gcc/Makefile.in @@ -2244,6 +2244,8 @@ version.o: $(REVISION) $(DATESTAMP) $(BASEVER) $(DEVPHASE) # lto-compress.o needs $(ZLIBINC) added to the include flags. CFLAGS-lto-compress.o += $(ZLIBINC) +CFLAGS-lto-streamer-in.o += -DTARGET_MACHINE=\"$(target_noncanonical)\" + bversion.h: s-bversion; @true s-bversion: BASE-VER echo "#define BUILDING_GCC_MAJOR `echo $(BASEVER_c) | sed -e 's/^\([0-9]*\).*$$/\1/'`" > bversion.h diff --git a/gcc/lto-streamer-in.c b/gcc/lto-streamer-in.c index 128d7640726..f59e8c5363d 100644 --- a/gcc/lto-streamer-in.c +++ b/gcc/lto-streamer-in.c @@ -1700,7 +1700,19 @@ lto_input_mode_table (struct lto_file_decl_data *file_data) } /* FALLTHRU */ default: - fatal_error (UNKNOWN_LOCATION, "unsupported mode %qs", mname); + /* For offloading-target compilions, this is a user-facing + message. See also target.def and machmode.def. */ + if (strcmp (mname, "XF") == 0) + fatal_error (UNKNOWN_LOCATION, + "%s - 80-bit floating-point numbers unsupported " + "(mode %qs)", TARGET_MACHINE, mname); + else if (strcmp (mname, "TF") == 0) + fatal_error (UNKNOWN_LOCATION, + "%s - 128-bit floating-point numbers unsupported " + "(mode %qs)", TARGET_MACHINE, mname); + else + fatal_error (UNKNOWN_LOCATION, "%s - unsupported mode %qs", + TARGET_MACHINE, mname); break; } } diff --git a/libgomp/testsuite/libgomp.c/pr86416-1.c b/libgomp/testsuite/libgomp.c/pr86416-1.c new file mode 100644 index 00000000000..c7a162d2f41 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/pr86416-1.c @@ -0,0 +1,22 @@ +/* { dg-do link } */ +/* { dg-require-effective-target large_long_double } */ + +/* PR middle-end/86416 */ +/* { dg-error "80-bit floating-point numbers unsupported .mode 'XF'." "" { target offload_device } 0 } */ +/* { dg-excess-errors "Follow-up errors from mkoffload and lto-wrapper" { target offload_device } } */ + +#include <stdlib.h> /* For abort. */ + +long double foo (long double x) +{ + #pragma omp target map(tofrom:x) + x *= 2.0; + return x; +} + +int main() +{ + long double v = foo (10.0q) - 20.0q; + if (v > 1.0e-5 || v < -1.0e-5) abort(); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/pr86416-2.c b/libgomp/testsuite/libgomp.c/pr86416-2.c new file mode 100644 index 00000000000..75e9cd773a3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/pr86416-2.c @@ -0,0 +1,22 @@ +/* { dg-do link { target __float128 } } */ +/* { dg-add-options __float128 } */ + +/* PR middle-end/86416 */ +/* { dg-error "128-bit floating-point numbers unsupported .mode 'TF'." "" { target offload_device } 0 } */ +/* { dg-excess-errors "Follow-up errors from mkoffload and lto-wrapper" { target offload_device } } */ + +#include <stdlib.h> /* For abort. */ + +__float128 foo(__float128 y) +{ + #pragma omp target map(tofrom: y) + y *= 4.0; + return y; +} + +int main() +{ + __float128 v = foo (5.0L) - 20.0L; + if (v > 1.0e-5 || v < -1.0e-5) abort(); + return 0; +}
As long as one compiles for a single target, the message is unlikely to appear. However, when compiling for offloading, the modes supported on the target 'host' and on the target 'device' can be different. In particular, 'long double' (when larger than double) and '__float128' might not be available on the device. This gives currently errors like the following (see PR, comment 0): > lto1: fatal error: unsupported mode TF > > compilation terminated. > mkoffload: fatal error: x86_64-pc-linux-gnu-accel-nvptx-none-gcc returned 1 exit status While the device target is hidden in 'x86_64-pc-linux-gnu-accel-nvptx-none-gcc', it might make more sense to add it more prominently. Additionally, an average user does not understand what 'TF' or 'XF' means. Solutions: (A) Add the target to the output (B) Add a better description for the error I did both in the attached patch, giving: lto1: fatal error: nvptx-none - 80-bit floating-point numbers unsupported (mode 'XF') lto1: fatal error: nvptx-none - 128-bit floating-point numbers unsupported (mode 'TF') * (A) should be fine, I think. * But I am not 100% happy with (B). I think as interim solution, it is acceptable as XF/TF are well defined and probably the most common problem. — Alternatively, one could only commit (A) or solve it more properly (how?). * If, e.g., 'long long' or 'integer(kind=16)' are used, only a generic message is printed. A message such as "'__float128' not supported" or "'real(kind=10)' not supported" is more helpful and supporting all modes and not cherry picking those two would be useful as well. The question is how to pass this information to lto-streamer-in.c; it is available as TYPE_NAME – and, with debugging turned on, this also gets passed on, but is also not be readily available in lto_input_mode_table. – Suggestions? Build on x86-64-gnu-linux and tested without offloading and with nvptx offloading. Tobias