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
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;
+}
