Richard Biener wrote:
On Fri, Sep 25, 2026 at 11:29 AM Tobias Burnus <[email protected]> wrote:
Richard Biener wrote:
Looks good, but how we get to pass == 3 is totally non-obvious ... why make it
conditional on pass == 3 only?
Maybe, we should add a comment to the code like:
--- a/gcc/lto-streamer-in.cc
+++ b/gcc/lto-streamer-in.cc
@@ -2141,5 +2141,6 @@ lto_input_mode_table (struct lto_file_decl_data
*file_data)
}
/* First search just the GET_CLASS_NARROWEST_MODE to wider modes,
- if not found, fallback to all modes. */
+ if not found, fallback to all modes. When found, pass is 3 (2 + 1),
+ when not found, pass is 2 ('< 2' check + 1). */
int pass;
for (pass = 0; pass < 2; pass++)
[...]
Thus, how about the patch plus adding a comment like above?
That would be nice, yes. But then the check could also be moved into
that 'else'?
You mean as follows?
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 | 25 +++++++++++++++++++++++--
libgomp/testsuite/libgomp.c/float16-target-1.c | 25 +++++++++++++++++++++++++
2 files changed, 48 insertions(+), 2 deletions(-)
diff --git a/gcc/lto-streamer-in.cc b/gcc/lto-streamer-in.cc
index 668a5d3065b..34c46b7da82 100644
--- a/gcc/lto-streamer-in.cc
+++ b/gcc/lto-streamer-in.cc
@@ -2136,15 +2136,16 @@ lto_input_mode_table (struct lto_file_decl_data *file_data)
real_fmt_name = bp_unpack_indexed_string (data_in, &bp,
&real_fmt_len);
break;
default:
break;
}
/* First search just the GET_CLASS_NARROWEST_MODE to wider modes,
- if not found, fallback to all modes. */
+ if not found, fallback to all modes. When found, pass is 3 (2 + 1),
+ when not found, pass is 2 ('< 2' check + 1). */
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) (mr + 1)
: mr = GET_MODE_WIDER_MODE (mr).else_void ())
@@ -2159,16 +2160,36 @@ lto_input_mode_table (struct lto_file_decl_data *file_data)
|| maybe_ne (GET_MODE_NUNITS (mr), nunits))
continue;
else if ((mclass == MODE_FLOAT || mclass == MODE_DECIMAL_FLOAT)
&& strcmp (REAL_MODE_FORMAT (mr)->name, real_fmt_name) != 0)
continue;
else
{
+ /* Found. */
table[m] = mr;
- pass = 2;
+ pass = 2; /* Will be 3 after the outer loop. */
+#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. */
+ 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 (mr)))
+ pass = 1; /* Such that pass == 2 after the outer loop. */
+ break;
+ default:
+ break;
+ }
+#endif
break;
}
unsigned int mname_len;
const char *mname = bp_unpack_indexed_string (data_in, &bp, &mname_len);
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;
+}
