https://gcc.gnu.org/g:221d38405f99e5c51f26e09759e79aee482611fa
commit r17-4657-g221d38405f99e5c51f26e09759e79aee482611fa Author: Tobias Burnus <[email protected]> Date: Fri Sep 25 17:10:55 2026 +0200 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. Diff: --- 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 668a5d3065b2..663952de7049 100644 --- a/gcc/lto-streamer-in.cc +++ b/gcc/lto-streamer-in.cc @@ -2140,7 +2140,8 @@ lto_input_mode_table (struct lto_file_decl_data *file_data) 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 @@ -2163,8 +2164,19 @@ lto_input_mode_table (struct lto_file_decl_data *file_data) 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; 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 000000000000..dfa12a69ad2b --- /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 000000000000..38d8dc701247 --- /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 000000000000..eb1975bf877c --- /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 000000000000..da39eaac1d9e --- /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 000000000000..ddcdf0cddca1 --- /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; +}
