gdb, testsuite: improve gdb.rocm cooperative group GWS tests

Refine the cooperative-group GWS tests for robustness and to make the
debugger behaviour under test more explicit:

  * coop-group-grid-sync.exp: use temporary breakpoints for the
    in-kernel locations (still pending, since the device code is
    loaded at dispatch time) and drop the redundant delete_breakpoints
    calls.

  * coop-group-grid-sync.exp: report UNSUPPORTED instead of FAIL when
    the inferior self-skips because no device advertises
    cooperativeLaunch.

  * coop-group-grid-sync.exp: have test_threads_in_coop_kernel check
    distinct per-wave blockIdx.x (per-wave register context) and
    per-lane threadIdx.x divergence (per-lane SIMT state), instead of
    only confirming that the backtrace names the kernel.

  * coop-group-multi-grid-sync.{cpp,exp}: read n_gpus from a marker
    line that is reached on every execution (no early return before
    it), so that when fewer than two cooperative-capable GPUs are
    available -- including when a parallel test run restricts the
    visible GPUs -- the test reports UNSUPPORTED rather than FAILing.

  * coop-group-{grid,multi-grid}-sync.{cpp,exp}: minor comment and
    GNU-style cleanups -- tab-align the in-kernel marker comment, keep
    hipLaunchCooperativeKernelMultiDevice on a single line, and
    clarify the Phase 2 data-dependency comment.

  * coop-group-{grid,multi-grid}-sync.{cpp,exp}: give the per-wave and
    per-lane value reads explicit, unique test names, and keep the
    "n-gpus-final" marker string unique so gdb_get_line_number
    resolves the intended line.

Tested on gfx942 with an in-tree build and the 7.14 nightly rocgdb,
both with all GPUs visible and with the visible set restricted to one.

Author: spatrang

gdb/testsuite/gdb.rocm/coop-group-grid-sync.cpp

@@ -40,8 +40,9 @@ namespace cg = cooperative_groups;
    in the GWS barrier.  Two workgroups of 64 threads each give 128
    threads in total.
 
-   N is exactly one slot per thread so Phase 2 has no intra-thread race
-   and the host-side expected values can be computed straightforwardly.  */
+   N is exactly one slot per thread so Phase 2 has no inter-thread
+   write conflict and the host-side expected values can be computed
+   straightforwardly.  */
 
 constexpr unsigned int group_size = 64;
 constexpr unsigned int num_groups = 2;

gdb/testsuite/gdb.rocm/coop-group-grid-sync.exp

@@ -55,10 +55,22 @@ proc_with_prefix test_break_around_grid_sync {} {
 	# Place a breakpoint before grid.sync ().  Use "allow-pending"
 	# because the kernel is loaded as device code at dispatch time.
 	gdb_breakpoint \
-	    [gdb_get_line_number "before-sync line"] allow-pending
-
-	# The kernel must hit the pre-sync breakpoint.
-	if {[gdb_continue_to_breakpoint "stop before grid.sync"] != 0} {
+	    [gdb_get_line_number "before-sync line"] allow-pending temporary
+
+	# Continue into the kernel.  If the program self-skips (no device
+	# advertises cooperativeLaunch) it exits cleanly; treat that as
+	# UNSUPPORTED rather than a FAIL.
+	set reached 0
+	gdb_test_multiple "continue" "stop before grid.sync" {
+	    -re -wrap "Temporary breakpoint $::decimal,\[^\r\n\]*coop_grid_sync_kernel .*" {
+		set reached 1
+		pass $gdb_test_name
+	    }
+	    -re -wrap "$::inferior_exited_re normally.*" {
+		unsupported "cooperative launch not available"
+	    }
+	}
+	if { !$reached } {
 	    return
 	}
 
@@ -90,27 +102,25 @@ proc_with_prefix test_break_around_grid_sync {} {
 	    ".*AMDGPU Dispatch\[^\r\n\]*coop_grid_sync_kernel.*" \
 	    "info dispatches lists cooperative dispatch"
 
-	# Replace the pre-sync breakpoint with one after grid.sync ().
-	# The after-sync breakpoint fires once all waves have crossed
-	# the GWS barrier, which proves GWS-protected code can be
-	# debugged across the barrier.
-	delete_breakpoints
+	# Place a breakpoint after grid.sync ().  It fires once all waves
+	# have crossed the GWS barrier, which proves GWS-protected code
+	# can be debugged across the barrier.
 	gdb_breakpoint \
-	    [gdb_get_line_number "after-sync line"] allow-pending
+	    [gdb_get_line_number "after-sync line"] allow-pending temporary
 
 	gdb_continue_to_breakpoint "stop after grid.sync"
 
 	# The kernel should be able to run to completion afterwards.
-	delete_breakpoints
 	gdb_continue_to_end "continue to program end" "continue" 1
     }
 }
 
-# Test: place a breakpoint inside the cooperative kernel, then iterate
-# over every stopped GPU wave and verify that each one has a backtrace
-# pointing into coop_grid_sync_kernel.  This exercises the wave
-# organization a debugger sees when several waves participate in a GWS
-# barrier.
+# Test: stop inside the cooperative kernel with several co-resident
+# waves parked at the barrier and check that gdb reports correct,
+# distinct per-wave and per-lane state.  This is the debugger
+# behaviour of interest: that gdb selects the right register context
+# per wave and the right SIMT state per lane, not merely that the
+# waves are in the kernel.
 
 proc_with_prefix test_threads_in_coop_kernel {} {
     clean_restart $::testfile
@@ -121,29 +131,56 @@ proc_with_prefix test_threads_in_coop_kernel {} {
 	}
 
 	gdb_breakpoint \
-	    [gdb_get_line_number "before-sync line"] allow-pending
-
-	if {[gdb_continue_to_breakpoint "stop in cooperative kernel"] != 0} {
+	    [gdb_get_line_number "before-sync line"] allow-pending temporary
+
+	# Continue into the kernel.  If the program self-skips (no device
+	# advertises cooperativeLaunch) it exits cleanly; treat that as
+	# UNSUPPORTED rather than a FAIL.
+	set reached 0
+	gdb_test_multiple "continue" "stop in cooperative kernel" {
+	    -re -wrap "Temporary breakpoint $::decimal,\[^\r\n\]*coop_grid_sync_kernel .*" {
+		set reached 1
+		pass $gdb_test_name
+	    }
+	    -re -wrap "$::inferior_exited_re normally.*" {
+		unsupported "cooperative launch not available"
+	    }
+	}
+	if { !$reached } {
 	    return
 	}
 
-	# Collect the list of AMDGPU waves and check that every one of
-	# them has a backtrace pointing into coop_grid_sync_kernel.
 	set waves [info_thread_get_wave_list]
 	gdb_assert {[llength $waves] >= 2} "multiple GPU waves listed"
 
+	# Per-wave register context: switching to each wave and reading
+	# blockIdx.x must yield that wave's owning workgroup, and across
+	# all waves we must observe more than one distinct workgroup.
+	set blocks {}
 	foreach wave $waves {
 	    with_test_prefix "wave $wave" {
 		gdb_test "thread $wave" "Switching to thread $wave.*" \
 		    "switch to wave $wave"
-		gdb_test "bt 1" \
-		    "#0\[^\r\n\]*coop_grid_sync_kernel\[^\r\n\]*" \
-		    "backtrace inside coop_grid_sync_kernel"
+		set bidx [get_integer_valueof "blockIdx.x" -1 "blockIdx.x"]
+		if {[lsearch -exact $blocks $bidx] == -1} {
+		    lappend blocks $bidx
+		}
 	    }
 	}
+	gdb_assert {[llength $blocks] >= 2} \
+	    "distinct workgroups observed across waves"
+
+	# Per-lane (SIMT) state: within a single wave, different lanes map
+	# to different threads, so threadIdx.x must differ between lanes.
+	gdb_test "thread [lindex $waves 0]" "Switching to thread .*" \
+	    "switch to first wave for lane inspection"
+	gdb_test "lane 0" ".*" "switch to lane 0"
+	set tid0 [get_integer_valueof "threadIdx.x" -1 "threadIdx.x at lane 0"]
+	gdb_test "lane 1" ".*" "switch to lane 1"
+	set tid1 [get_integer_valueof "threadIdx.x" -1 "threadIdx.x at lane 1"]
+	gdb_assert {$tid0 != $tid1} "distinct threadIdx.x across lanes"
 
 	# Resume to completion.
-	delete_breakpoints
 	gdb_continue_to_end "continue to program end" "continue" 1
     }
 }

gdb/testsuite/gdb.rocm/coop-group-multi-grid-sync.cpp

@@ -62,7 +62,7 @@ coop_multi_grid_sync_kernel (int *data, unsigned int n_elements,
 
   /* Phase 1: each thread writes data into its slot.  */
   for (unsigned int i = tid; i < n_elements; i += stride)
-    data[i] = (int) ((grid_rank + 1) * (i + 1));   /* before-grid-sync line */
+    data[i] = (int) ((grid_rank + 1) * (i + 1));	/* before-grid-sync line */
 
   /* Intra-device grid sync (GWS).  */
   grid.sync ();			       /* grid-sync line */
@@ -97,17 +97,14 @@ main ()
 {
   int n_devices = 0;
   CHECK (hipGetDeviceCount (&n_devices));
-  if (n_devices < N_USED_GPUS)
-    {
-      printf ("Multi-device cooperative test needs >= %d GPUs"
-	      " (found %d), skipping.\n", N_USED_GPUS, n_devices);
-      return 0;
-    }
 
   /* Pick the first N_USED_GPUS devices that support
      cooperativeMultiDeviceLaunch.  In a mixed-architecture system not
      every GPU has GWS support, so we ignore unsupported ones rather
-     than failing the whole test.  */
+     than failing the whole test.  This loop runs unconditionally (it
+     is bounded by n_devices) so the marker below is always reached,
+     even when too few devices are available; the companion .exp reads
+     n_gpus there to decide whether to skip.  */
   int selected[N_USED_GPUS];
   int n_gpus = 0;
   for (int id = 0; id < n_devices && n_gpus < N_USED_GPUS; id++)
@@ -120,15 +117,14 @@ main ()
 	  n_gpus++;
 	}
     }
-  if (n_gpus < N_USED_GPUS)
+  if (n_gpus < N_USED_GPUS)	/* n-gpus-final line */
     {
-      printf ("Fewer than %d devices in 0..%d support cooperative"
+      printf ("Fewer than %d of the %d HIP device(s) support cooperative"
 	      " multi-device launch, skipping.\n",
-	      N_USED_GPUS, n_devices - 1);
+	      N_USED_GPUS, n_devices);
       return 0;
     }
 
-  /* n-gpus-final line.  */
   int *data_d[N_USED_GPUS] = {};
   hipStream_t stream[N_USED_GPUS] = {};
 

gdb/testsuite/gdb.rocm/coop-group-multi-grid-sync.exp

@@ -26,8 +26,8 @@
 #   - read the number of participating GPUs from the inferior,
 #   - hit a kernel breakpoint inside the cooperative dispatch on
 #     *every* participating GPU (matched via the per-device
-#     "Breakpoint <id>.<instance>" subscript that hipLaunchCooperative
-#     KernelMultiDevice produces), and
+#     "Breakpoint <id>.<instance>" subscript that
+#     hipLaunchCooperativeKernelMultiDevice produces), and
 #   - resume the cross-device dispatch through both the intra-device
 #     this_grid ().sync () and the cross-device this_multi_grid ().sync ()
 #     to normal program completion.
@@ -63,34 +63,19 @@ proc_with_prefix test_break_in_multi_coop_kernel {} {
 	    return
 	}
 
-	# Advance to the point where the inferior has finalized "n_gpus"
-	# (the count of GPUs it will actually use for the cooperative
-	# dispatch) so we can read it from gdb.  The require gates above
-	# cover architecture-level support; per-device runtime gating
-	# (>= 2 devices that advertise cooperativeMultiDeviceLaunch) is
-	# done inside main () in the .cpp, which self-skips with a clean
-	# exit when those preconditions are not met.  Treat that clean
-	# exit as UNSUPPORTED rather than a FAIL.
-	gdb_breakpoint \
-	    [gdb_get_line_number "n-gpus-final line"] allow-pending
-
-	set eligible 0
-	gdb_test_multiple "continue" "advance to n-gpus-final" {
-	    -re "hit Breakpoint $::decimal,\[^\r\n\]*main \\(\\)\[^\r\n\]*\r\n\[^\r\n\]*\r\n$::gdb_prompt " {
-		set eligible 1
-		pass $gdb_test_name
-	    }
-	    -re "\\\[Inferior 1 \[^\r\n\]* exited normally\\\]\[^\r\n\]*\r\n$::gdb_prompt " {
-		unsupported "multi-device cooperative launch not available"
-	    }
-	}
-	if { !$eligible } {
-	    return
-	}
+	# Break where the inferior has finalized "n_gpus" (the count of
+	# GPUs it will use for the cooperative dispatch).  This line runs
+	# on every execution, before the inferior's own eligibility
+	# check, so reading "n_gpus" from gdb lets us decide here whether
+	# to proceed.  The require gates above cover architecture-level
+	# support; per-device runtime gating (>= 2 devices that advertise
+	# cooperativeMultiDeviceLaunch) can only be known at this point.
+	gdb_breakpoint [gdb_get_line_number "n-gpus-final line"] temporary
+	gdb_continue_to_breakpoint "advance to n-gpus-final"
 
 	set n_gpus [get_integer_valueof "n_gpus" 0]
-	gdb_assert {$n_gpus >= 2} "n_gpus is at least 2"
 	if { $n_gpus < 2 } {
+	    unsupported "fewer than two cooperative-capable GPUs"
 	    return
 	}
 
@@ -107,12 +92,11 @@ proc_with_prefix test_break_in_multi_coop_kernel {} {
 	    }
 	}
 
-	# In non-stop mode, hipLaunchCooperativeKernelMultiDevice
-	# produces one child breakpoint instance per participating GPU
-	# ("Breakpoint <id>.<inst>").  Collect distinct <inst> values
-	# until we have observed a stop on every GPU; only then is it
-	# safe to delete the breakpoint and let the dispatch run
-	# through both grid syncs.
+	# Check that gdb resolves the single source breakpoint to one
+	# child location per participating GPU ("Breakpoint <id>.<inst>")
+	# and, in non-stop mode, reports each device-side stop
+	# independently.  Collect distinct <inst> values until a stop has
+	# been seen for every GPU.
 	array unset gpu_seen
 	set distinct_gpus 0
 	gdb_test_multiple "" "breakpoint hit on every participating GPU" {