That's with OpenMP offloading to either nvptx or amdgcn.
Some benchmark code nicely compiled on the host - and then
the device lto1 compiler failed with:
lto1: internal compiler error: tree check:
expected class ‘type’, have ‘exceptional’ (error_mark)
in lto_fixup_prevailing_type, at lto/lto-common.cc:1006
It turned out that the code in question used _Float16.
For _Float128, no machine mode is found and a
nice error message is shown. But not so for _Float16.
The reason, it turned out: "HF" is actually available as
machine mode – but the targetm.scalar_mode_supported_p
hook then simply reports it as unsupported, effectively
disabling it.
(In the amdgcn's hook, HF is is commented out; in nvptx's,
it only available with -mexperimental [if additionally
-march=sm_53 or higher].)
* * *
When using the nvptx/amdgcn C compiler, _Float16 was rejected
during parsing; however, only with the attached patch there
now also a proper error message for offloaded code:
lto1: fatal error: nvptx-none - 16-bit-precision floating-point numbers
unsupported (mode ‘HF’)
lto1: fatal error: amdgcn-amdhsa - 16-bit-precision floating-point numbers
unsupported (mode ‘HF’)
OK for mainline?
Tobias
lto_input_mode_table: on offload targets, check whether mode is actually available [PR127589]
For an offload-target compiler, the modes in the code match the modes on the
host compiler - and not all modes might be available on the device side.
There was already a check whether a mode was not existing at all, e.g.
_Float128 are (currently?) unsupported on nvptx and amdgcn offload devices.
However, for _Float16 ("HF") is more interesting: Both mentioned offload
devices have support for it, but it is not enabled at runtime (amdgcn)
or only sometimes (nvptx: only with -mexperimental for sm_53+ devices).
Thus, an additional runtime check has been added for ACCEL_COMPILER, only
to call the hook targetm.scalar_mode_supported_p to check whether the
mode is actually available. And if not, an error message is printed.
PR lto/127589
gcc/ChangeLog:
* lto-streamer-in.cc (lto_input_mode_table): With ACCEL_COMPILER,
check whether found float and int modes are actually enabled/available
at runtime.
libgomp/ChangeLog:
* testsuite/libgomp.c/float16-target-1.c: New test.
gcc/lto-streamer-in.cc | 21 +++++++++++++++++++++
libgomp/testsuite/libgomp.c/float16-target-1.c | 25 +++++++++++++++++++++++++
2 files changed, 46 insertions(+)
diff --git a/gcc/lto-streamer-in.cc b/gcc/lto-streamer-in.cc
index 668a5d3065b..b5a800a8d08 100644
--- a/gcc/lto-streamer-in.cc
+++ b/gcc/lto-streamer-in.cc
@@ -2169,6 +2169,27 @@ lto_input_mode_table (struct lto_file_decl_data *file_data)
}
unsigned int mname_len;
const char *mname = bp_unpack_indexed_string (data_in, &bp, &mname_len);
+
+#ifdef ACCEL_COMPILER
+ /* A mode used by the host code - and generally supported by the
+ offloading-target compiler might not actually be available, depending
+ e.g. on -march or other other commandline options. Hence, check again
+ and, if failing, diagnose it below. */
+ if (pass == 3)
+ switch (mclass)
+ {
+ case MODE_FLOAT:
+ case MODE_DECIMAL_FLOAT:
+ case MODE_COMPLEX_FLOAT:
+ case MODE_INT:
+ if (!targetm.scalar_mode_supported_p (scalar_mode::from_int (table[m])))
+ pass = 2;
+ break;
+ default:
+ break;
+ }
+#endif
+
if (pass == 2)
{
switch (mclass)
diff --git a/libgomp/testsuite/libgomp.c/float16-target-1.c b/libgomp/testsuite/libgomp.c/float16-target-1.c
new file mode 100644
index 00000000000..da39eaac1d9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/float16-target-1.c
@@ -0,0 +1,25 @@
+/* { dg-do link { target float16 } } */
+/* { dg-add-options float16 } */
+
+/* PR lto/127589 */
+
+/* Both NVPTX and AMDGCN define HF - and support it in principle (nvptx: sm_XX > 53)
+ However, both have currently disabled it at runtime (target hook)
+ [For Nvpts, -mexperimental can actually be used to enable it.]
+
+ Hence: Check for the following - otherwise, it would ICE. */
+
+/* { dg-error "bit-precision floating-point numbers unsupported .mode '.F'." "" { target { offload_target_nvptx || offload_target_amdgcn } } 0 } */
+/* { dg-excess-errors "Follow-up errors from mkoffload and lto-wrapper" { target { offload_target_nvptx || offload_target_amdgcn } } } */
+
+int f(_Float16 x) {
+ return x > 0.0;
+}
+
+int main() {
+ int r;
+ _Float16 y = (_Float16)0.0;
+ #pragma omp target map(from: r), map(to: y)
+ r = f(y);
+ return r;
+}
