Richard Biener wrote:
On Fri, Sep 25, 2026 at 1:56 PM Tobias Burnus <[email protected]> wrote:
Or rather like this – with additional tests, esp. for _Complex
and _Decimal.
Yes. Though you can possibly simplify that switch () by using
scalar_mode::includes_p (mr).
Which would be this variant ...
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/cmplxint-target-1.c: New test.
* testsuite/libgomp.c/cmplxint128-target-1.c: New test.
* testsuite/libgomp.c/dfp-target-1.c: New test.
* testsuite/libgomp.c/float16-target-1.c: New test.
* testsuite/libgomp.c/float16-target-2.c: New test.
gcc/lto-streamer-in.cc | 16 ++++++++++++--
libgomp/testsuite/libgomp.c/cmplxint-target-1.c | 21 ++++++++++++++++++
libgomp/testsuite/libgomp.c/cmplxint128-target-1.c | 21 ++++++++++++++++++
libgomp/testsuite/libgomp.c/dfp-target-1.c | 18 ++++++++++++++++
libgomp/testsuite/libgomp.c/float16-target-1.c | 25 ++++++++++++++++++++++
libgomp/testsuite/libgomp.c/float16-target-2.c | 25 ++++++++++++++++++++++
6 files changed, 124 insertions(+), 2 deletions(-)
diff --git a/gcc/lto-streamer-in.cc b/gcc/lto-streamer-in.cc
index 668a5d3065b..663952de704 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,27 @@ 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. */
+ if (scalar_mode::includes_p (mr)
+ && !targetm.scalar_mode_supported_p (
+ scalar_mode::from_int (mr)))
+ pass = 1; /* Such that pass == 2 after the outer loop. */
+#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/cmplxint-target-1.c b/libgomp/testsuite/libgomp.c/cmplxint-target-1.c
new file mode 100644
index 00000000000..dfa12a69ad2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/cmplxint-target-1.c
@@ -0,0 +1,21 @@
+/* { dg-do link } */
+
+/* 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. */
+
+int f(_Complex int x) {
+ return __real__ x > 0;
+}
+
+int main() {
+ int r;
+ _Complex int y = (_Complex int)0;
+ #pragma omp target map(from: r), map(to: y)
+ r = f(y);
+ return r;
+}
diff --git a/libgomp/testsuite/libgomp.c/cmplxint128-target-1.c b/libgomp/testsuite/libgomp.c/cmplxint128-target-1.c
new file mode 100644
index 00000000000..38d8dc70124
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/cmplxint128-target-1.c
@@ -0,0 +1,21 @@
+/* { dg-do link { target int128 } } */
+
+/* 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. */
+
+int f(_Complex __int128 x) {
+ return __real__ x > 0;
+}
+
+int main() {
+ int r;
+ _Complex __int128 y = (_Complex __int128)0;
+ #pragma omp target map(from: r), map(to: y)
+ r = f(y);
+ return r;
+}
diff --git a/libgomp/testsuite/libgomp.c/dfp-target-1.c b/libgomp/testsuite/libgomp.c/dfp-target-1.c
new file mode 100644
index 00000000000..eb1975bf877
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/dfp-target-1.c
@@ -0,0 +1,18 @@
+/* { dg-do link { target dfp } } */
+
+/* PR lto/127589 */
+
+/* { dg-error "32-bit-precision decimal floating-point numbers unsupported .mode 'SD'." "" { 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(_Decimal32 x) {
+ return x > (_Decimal32)0.0;
+}
+
+int main() {
+ int r;
+ _Decimal32 y = (_Decimal32)0.0;
+ #pragma omp target map(from: r), map(to: y)
+ r = f(y);
+ return r;
+}
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;
+}
diff --git a/libgomp/testsuite/libgomp.c/float16-target-2.c b/libgomp/testsuite/libgomp.c/float16-target-2.c
new file mode 100644
index 00000000000..ddcdf0cddca
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/float16-target-2.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(_Complex _Float16 x) {
+ return __real__ x > (_Float16)0.0;
+}
+
+int main() {
+ int r;
+ _Complex _Float16 y = (_Complex _Float16)0.0;
+ #pragma omp target map(from: r), map(to: y)
+ r = f(y);
+ return r;
+}
