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.
---
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 ece41c59bbb..92de6fb1b64 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 2fc811d91c6..076b775560f 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
;
--
2.51.0
