[PATCH 1/2] gdb/testsuite: For ROCm/HIP, don't rely on rocm_agent_enumerator

The rocm_agent_enumerator program does not exist on all HIP supported
platforms (such as e.g. the Windows HIP SDK).

Fix this by compiling and running our own replacement HIP program that
lists AMD GPU devices, instead of calling rocm_agent_enumerator.

With this, we no longer need to check if we have a working HIP
compiler available (by compiling a similar HIP program), as compiling
the enumerator program achieves the same goal.  In turn this means
that the cost of having this replacement is essentially zero.  Same
number of external process invocations, and essentially the same
number of lines of code.  And we get to drop one dependency.

We just need to be careful to not call hcc_amdgpu_targets when
compiling the GPU device enumerator program, otherwise we hit infinite
recursion.  That is handled by passing a new hip_no_offload_arch
option to gdb_compile.

Both Linux and Windows ROCm builds nowadays ship with an alternative
amdgpu-arch program (also installed as offload-arch) that we could
use, but as explained above, having our own replacement has basically
no cost, so I'm not proposing using it.

Note also that the hip_no_offload_arch flag will be used in other
places, in future patches.

Change-Id: I45262f2fe24455ade6075eded2cd7df03979c6cc
---
 gdb/testsuite/lib/gdb.exp  |  4 +-
 gdb/testsuite/lib/rocm.exp | 93 +++++++++++++++++++-------------------
 2 files changed, 50 insertions(+), 47 deletions(-)

diff --git a/gdb/testsuite/lib/gdb.exp b/gdb/testsuite/lib/gdb.exp
index d4b8c9c24dc..9bb5a1446e6 100644
--- a/gdb/testsuite/lib/gdb.exp
+++ b/gdb/testsuite/lib/gdb.exp
@@ -6554,6 +6554,7 @@ proc gdb_windows_manifest_obj {} {
 #   - column-info/no-column-info: Enable/Disable generation of column table
 #     information.
 #   - dwarf5: Force compilation with dwarf-5 debug information.
+#   - hip_no_offload_arch: Do not pass --offload-arch to the HIP compiler.
 #
 # And here are some of the not too obscure options understood by DejaGnu that
 # influence the compilation:
@@ -6974,7 +6975,8 @@ proc gdb_compile {source dest type options} {
 	# amdclang++ requires explicit --offload-arch.  Explicitly
 	# pass one --offload-arch for each available device.  But
 	# don't do it if the testcase explicitly used --offload-arch.
-	if {[lsearch -regexp $options "--offload-arch="] == -1} {
+	if {[lsearch -exact $options hip_no_offload_arch] == -1
+	    && [lsearch -regexp $options "--offload-arch="] == -1} {
 	    foreach gpu_target [hcc_amdgpu_targets] {
 		lappend new_options "early_flags=--offload-arch=$gpu_target"
 	    }
diff --git a/gdb/testsuite/lib/rocm.exp b/gdb/testsuite/lib/rocm.exp
index f8a3b233fae..025fa63da6f 100644
--- a/gdb/testsuite/lib/rocm.exp
+++ b/gdb/testsuite/lib/rocm.exp
@@ -46,38 +46,58 @@ proc log_host_exec { cmd } {
 #
 # Return a list of GPU devices that do exist on the system.
 # The list will be empty when there's no GPU or the execution
-# of rocm_agent_enumerator does not succeed.  It is up to the
-# caller of this procedure that what should happen when an empty
+# of the enumerator program does not succeed.  It is up to the
+# caller of this procedure what should happen when an empty
 # list is returned.
 
 gdb_caching_proc find_amdgpu_devices {} {
-    global rocm_path
-    set hip_gpu_devices [list]
-    set enumerator "rocm_agent_enumerator"
-    set targets ""
+    # Compile and run a simple custom HIP program that lists all GPU
+    # devices, one device per line.  We don't rely on
+    # rocm_agent_enumerator because that does not exist on all
+    # supported platforms.
+    #
+    # Compile without --offload-arch (which is fine because this is a
+    # host-only program), because otherwise gdb_compile would call
+    # hcc_amdgpu_targets to know which --offload-arch flags to pass to
+    # the HIP compiler, and we'd end up here again, resulting in
+    # infinite recursion.
+    set options {hip hip_no_offload_arch}
+    if {![gdb_simple_compile device_enumerator {
+	    #include <hip/hip_runtime.h>
+	    #include <stdio.h>
+	    #include <string.h>
 
-    # Try the PATH first
-    set result [log_host_exec "$enumerator"]
-    if {[lindex $result 0] == 0} {
-	set targets [lindex $result 1]
-    } else {
-	# Now try the ROCM_PATH
-	set result [log_host_exec "$rocm_path/bin/$enumerator"]
-	if {[lindex $result 0] == 0} {
-	    set targets [lindex $result 1]
-	}
+	    int
+	    main ()
+	    {
+	      int device_count;
+	      if (hipGetDeviceCount (&device_count) == hipSuccess)
+		for (int i = 0; i < device_count; i++)
+		  {
+		    hipDeviceProp_t props;
+		    if (hipGetDeviceProperties (&props, i) == hipSuccess)
+		      {
+			/* Strip out the supported features list,
+			   like gfx90a:sramecc+:xnack-.  */
+			char *colon = strstr (props.gcnArchName, ":");
+			if (colon != nullptr)
+			  *colon = '\0';
+			printf ("%s\n", props.gcnArchName);
+		      }
+		  }
+	    }
+	} executable $options]} {
+	return {}
     }
 
-    if {$targets != ""} {
-	foreach dev $targets {
-	    # Ignore the 'gfx000' device which identifies the host.
-	    if {$dev != "gfx000"} {
-		lappend hip_gpu_devices $dev
-	    }
-	}
+    set result [log_host_exec "$obj"]
+    if {[lindex $result 0] == 0} {
+	set targets [lindex $result 1]
+	# Convert newline-separated string to a list.
+	return [list {*}$targets]
     }
 
-    return $hip_gpu_devices
+    return {}
 }
 
 # Get the list of unique GPU targets to compile for.
@@ -128,33 +148,14 @@ gdb_caching_proc allow_hip_tests {} {
 	return {0 "amd-dbgapi not supported"}
     }
 
-    # Check if there's any GPU device to run the tests on.
+    # Check if there's any GPU device to run the tests on.  If this
+    # works, then we also know we have a working HIP compiler
+    # available.
     set devices [find_amdgpu_devices]
     if {[llength $devices] == 0} {
 	return {0 "no suitable amdgpu targets found"}
     }
 
-    # Check if we have a working hipcc compiler available.
-    # TARGETS won't be empty, because there's at least one GPU device.
-    set targets [hcc_amdgpu_targets]
-    set flags [list hip additional_flags=--offload-arch=[join $targets ","]]
-    if {![gdb_simple_compile hipprobe {
-	    #include <hip/hip_runtime.h>
-	    __global__ void
-	    kern () {}
-
-	    int
-	    main ()
-	    {
-		kern<<<1, 1>>> ();
-		if (hipDeviceSynchronize () != hipSuccess)
-		  return -1;
-		return 0;
-	    }
-	} executable $flags]} {
-	return {0 "failed to compile hip program"}
-    }
-
     return 1
 }
 
-- 
2.54.0