https://gcc.gnu.org/g:1d1d12da6ddc0030e3038363d46fdd23f8cd068f
commit r16-5890-g1d1d12da6ddc0030e3038363d46fdd23f8cd068f Author: Andrew Stubbs <[email protected]> Date: Wed Dec 3 15:21:30 2025 +0000 amdgcn, libgomp: improve generic device errors Switching to use "generic" ISA variants has changed the error modes a bit. This patch changes the runtime so that it doesn't say to use the device-specific -march option when the real problem is not the ISA (it'll be a mismatched xnack setting, probably). Additionally, the testsuite effective target check needs to see if the xnack mode is accepted by the runtime, as well as the compiler. libgomp/ChangeLog: * plugin/plugin-gcn.c (generic_isa_code): New function. (isa_matches_agent): Use generic ISA details to help select an error message on ISA mismatch. * testsuite/lib/libgomp.exp (check_effective_target_offload_target_amdgcn_with_xnack): Use a runtime check. Diff: --- libgomp/plugin/plugin-gcn.c | 20 +++++++++++++++++++- libgomp/testsuite/lib/libgomp.exp | 2 +- 2 files changed, 20 insertions(+), 2 deletions(-) diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c index ece41c59bbb5..92de6fb1b64e 100644 --- a/libgomp/plugin/plugin-gcn.c +++ b/libgomp/plugin/plugin-gcn.c @@ -1776,6 +1776,22 @@ isa_code(const char *isa) { return EF_AMDGPU_MACH_UNSUPPORTED; } +/* Returns the code which is used in the GCN object code to identify the + generic ISA that corresponds to a specific ISA. */ + +static gcn_isa +generic_isa_code (int isa) { + switch(isa) + { +#define EF_AMDGPU_MACH_AMDGCN_NONE 0 +#define GCN_DEVICE(name, NAME, ELF, GCCISA, XNACK, SRAM, WAVE64, CUMODE, \ + VGPRS, CO, ARCH, GENERIC_ISA, ...) \ + case ELF: return EF_AMDGPU_MACH_AMDGCN_ ## GENERIC_ISA; +#include "../../gcc/config/gcn/gcn-devices.def" + } + return 0; +} + /* CDNA2 devices have twice as many VGPRs compared to older devices. */ static int @@ -2551,7 +2567,9 @@ isa_matches_agent (struct agent_info *agent, Elf64_Ehdr *image, "Consider using ROCR_VISIBLE_DEVICES to disable incompatible " "devices or run with LOADER_ENABLE_LOGGING=1 for more details.", device_isa_s, agent_isa_s, agent->device_id); - else if (strcmp (device_isa_s, agent_isa_s) == 0) + else if (strcmp (device_isa_s, agent_isa_s) == 0 + || (elf_gcn_isa_is_generic (image) + && generic_isa_code (agent->device_isa) == isa_field)) snprintf (msg, sizeof msg, "GCN code object features do not match for an unknown reason " "(device %d).\n" diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp index 2fc811d91c64..076b775560fa 100644 --- a/libgomp/testsuite/lib/libgomp.exp +++ b/libgomp/testsuite/lib/libgomp.exp @@ -758,7 +758,7 @@ proc check_effective_target_omp_managedmem { } { proc check_effective_target_offload_target_amdgcn_with_xnack { } { if { [libgomp_check_effective_target_offload_target "amdgcn"] } { - return [check_no_compiler_messages amd_xnack_ executable { + return [check_runtime amd_xnack_ { int main () { #pragma omp target ;
