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
2019-12-18 Tobias Burnus <tob...@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.
gcc/Makefile.in | 2 ++
gcc/lto-streamer-in.c | 26 +++++++++++++++++++++++++-
libgomp/testsuite/libgomp.c/pr86416-1.c | 22 ++++++++++++++++++++++
libgomp/testsuite/libgomp.c/pr86416-2.c | 22 ++++++++++++++++++++++
4 files changed, 71 insertions(+), 1 deletion(-)
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 675e1a7a153..f49f38df647 100644
--- a/gcc/lto-streamer-in.c
+++ b/gcc/lto-streamer-in.c
@@ -1698,7 +1698,31 @@ lto_input_mode_table (struct lto_file_decl_data *file_data)
}
/* FALLTHRU */
default:
- fatal_error (UNKNOWN_LOCATION, "unsupported mode %qs", mname);
+ /* This is only used for offloading-target compilations and
+ is a user-facing error. Give a better error message for
+ the common modes; see also mode-classes.def. */
+ if (mclass == MODE_FLOAT)
+ fatal_error (UNKNOWN_LOCATION,
+ "%s - %u-bit-precision floating-point numbers "
+ "unsupported (mode %qs)", TARGET_MACHINE,
+ prec.to_constant (), mname);
+ else if (mclass == MODE_DECIMAL_FLOAT)
+ fatal_error (UNKNOWN_LOCATION,
+ "%s - %u-bit-precision decimal floating-point "
+ "numbers unsupported (mode %qs)", TARGET_MACHINE,
+ prec.to_constant (), mname);
+ else if (mclass == MODE_COMPLEX_FLOAT)
+ fatal_error (UNKNOWN_LOCATION,
+ "%s - %u-bit-precision complex floating-point "
+ "numbers unsupported (mode %qs)", TARGET_MACHINE,
+ prec.to_constant (), mname);
+ else if (mclass == MODE_INT)
+ fatal_error (UNKNOWN_LOCATION,
+ "%s - %u-bit integer numbers unsupported (mode "
+ "%qs)", TARGET_MACHINE, prec.to_constant (), 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..4ab523d2310
--- /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 "bit-precision floating-point numbers unsupported .mode '.F'." "" { 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..f104da78029
--- /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 "bit-precision floating-point numbers unsupported .mode '.F'." "" { 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;
+}
