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 --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c
index ece41c59bbb50300a89747ffc3a2b9980a451d6e..92de6fb1b64e30019d1ff4d983ea4e340640568d 100644 (file)
--- 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 2fc811d91c64b28b8d7c9860fc126d0d7657211f..076b775560fa1c5aa4fe3dedab0974ead3d46437 100644 (file)
--- 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
               ;