Add dejagnu tests for cooperative group GWS debugging by spatrang · Pull Request #116 · ROCm/ROCgdb

spatrang · 2026-05-07T13:08:08Z

Summary

Add dejagnu coverage for debugging AMD GPU cooperative-group kernels —
i.e. kernels launched via hipLaunchCooperativeKernel /
hipLaunchCooperativeKernelMultiDevice that synchronize at the grid /
multi-grid level. On AMD GPUs these synchronization primitives are
implemented in hardware via Global Wave Sync (GWS), and they have a
distinct wave/scheduling model that has historically only been covered by
out-of-tree tests. This PR brings that coverage into the dejagnu testsuite
so it runs as part of the regular ROCgdb regression suite.

Tests added

File	Scenario
`gdb.rocm/coop-group-grid-sync.{cpp,exp}`	Single-device cooperative kernel using `cooperative_groups::this_grid().sync()` (intra-device GWS), launched via `hipLaunchCooperativeKernel`.
`gdb.rocm/coop-group-multi-grid-sync.{cpp,exp}`	Multi-device cooperative kernel using both `this_grid().sync()` and `cooperative_groups::this_multi_grid().sync()` (intra + cross-device GWS), launched via `hipLaunchCooperativeKernelMultiDevice`.

What gets verified

coop-group-grid-sync.exp — two sub-tests:

test_break_around_grid_sync
- Hit a breakpoint before grid.sync() inside a cooperative dispatch.
- Confirm multiple AMDGPU Wave threads are stopped (waves participating
  in the GWS barrier).
- Confirm info dispatches lists the cooperative dispatch.
- Move the breakpoint to after grid.sync() and continue: it must
  fire (proves GWS-protected code can be debugged across the barrier).
- Continue to clean program exit.
test_threads_in_coop_kernel
- For every AMDGPU Wave parked inside the kernel, switch to it and
  confirm bt 1 reports a frame inside coop_grid_sync_kernel.

coop-group-multi-grid-sync.exp — runs in non-stop mode:

After continue -a &, confirm a kernel-side breakpoint fires inside
coop_multi_grid_sync_kernel. Per-GPU child breakpoint instances
(Breakpoint X.Y) are observed for every participating GPU.
Continue all threads to program exit, which only succeeds if both
this_grid().sync() and this_multi_grid().sync() release correctly
under the debugger.

The host-side post-conditions in the .cpp programs additionally validate
the cooperative semantics numerically (cross-workgroup data dependency for
the single-device case, cross-device sum aggregation for the multi-device
case), so any regression in GWS behavior under the debugger turns into a
test failure rather than a silent miscompare.

Skip / unsupported handling

The tests degrade cleanly on systems that cannot run them:

Single-device test: queries cooperativeLaunch; if unsupported the
program prints a recognizable message and exits, and the .exp marks
the test UNSUPPORTED.
Multi-device test: requires >= 2 GPUs and
cooperativeMultiDeviceLaunch on every device. It is also gated by
the existing hip_devices_support_debug_multi_process requirement.
Any of those missing → UNSUPPORTED.

No new dejagnu helpers are required; both .exp files use existing
infrastructure in lib/rocm.exp.

Out of scope / follow-ups

Intentionally left out of this PR; happy to extend if reviewers ask:

Stepping (next / step / stepi) across grid.sync() /
mgrid.sync() boundaries.
Conditional breakpoints inside cooperative kernels.
lane apply / per-lane register inspection while waves are parked at
the GWS barrier.
Watchpoints on cooperative shared buffers.

Copilot

Pull request overview

Adds new ROCm dejagnu coverage to exercise ROCgdb debugging of cooperative-group HIP kernels that synchronize via GWS, covering both single-device this_grid().sync() and multi-device this_multi_grid().sync() scenarios.

Changes:

Introduces a single-device cooperative-kernel test that breaks before/after grid.sync() and validates waves/dispatch visibility.
Introduces a multi-device cooperative-kernel non-stop test that breaks inside a multi-grid kernel and runs through grid + multi-grid barriers to completion.
Adds two HIP C++ test programs that implement the cooperative-group synchronization patterns and validate results on the host side.

Reviewed changes

Copilot reviewed 4 out of 4 changed files in this pull request and generated 2 comments.

File	Description
gdb/testsuite/gdb.rocm/coop-group-grid-sync.exp	DejaGnu test for single-device cooperative kernel debugging around `this_grid().sync()`.
gdb/testsuite/gdb.rocm/coop-group-grid-sync.cpp	HIP program implementing single-device cooperative `grid.sync()` and host-side validation.
gdb/testsuite/gdb.rocm/coop-group-multi-grid-sync.exp	DejaGnu non-stop test for multi-device cooperative kernel debugging through `this_grid().sync()` + `this_multi_grid().sync()`.
gdb/testsuite/gdb.rocm/coop-group-multi-grid-sync.cpp	HIP program implementing multi-device cooperative launch with cross-device aggregation and validation.

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

lancesix

Hi,
Thanks a lot for this, this is a great starting point.

My main concern for now is gfx110x. We do not support debugging cooperative group on those (documented limitation), the testcase should look for them to not FAIL. This is known that the test will not pass even if the arch do support GWS.

I have added a couple of small comments, I'll get back to a more detailed review after the gfx11 concern has been addressed.

lancesix · 2026-05-07T13:44:48Z

+
+constexpr unsigned int group_size = 64;
+constexpr unsigned int num_groups = 2;
+constexpr int total_threads = group_size * num_groups;


Why are group_size and num_groups unsigned, but not total_threads?

lancesix · 2026-05-07T13:52:43Z

+	}
+	-re "\\\[Inferior 1 \[^\r\n\]* exited normally\\\]\[^\r\n\]*\r\n$::gdb_prompt " {
+	    # Program skipped without our diagnostic line, but still
+	    # exited cleanly.  Treat as unsupported.


Which case is this covering? Shouldn't there be an XFAIL if the environment is broken to a point where we do not see the message?

lancesix · 2026-05-07T13:53:43Z

+
+	# The kernel must hit the pre-sync breakpoint.
+	if {![continue_to_kernel_or_skip "stop before grid.sync"]} {
+	    return


We should report UNSUPPORTED here.

lancesix · 2026-05-07T13:54:38Z

+	# architectures).
+	set waves [info_thread_get_wave_list]
+	gdb_assert {[llength $waves] >= 2} \
+	    "at least two AMDGPU waves are present pre-sync"


Can there be a race where, when running on wave32, we just have the 2 waves of a single workgroup, but not the waves of the other?

lancesix · 2026-05-07T13:55:56Z

+	    [gdb_get_line_number "before-sync line"] allow-pending
+
+	if {![continue_to_kernel_or_skip "stop in cooperative kernel"]} {
+	    return


Report UNSUPPORTED here.

lancesix · 2026-05-07T13:57:05Z

+   hipLaunchCooperativeKernelMultiDevice.  The kernel uses both
+   cooperative_groups::this_grid ().sync () (intra-device GWS) and
+   cooperative_groups::this_multi_grid ().sync () (cross-device sync),
+   matching the historical CQE GWS test exercised on multiple ASICs.


I would probably say "historical out of tree" or something, CQE does not really mean much outside of AMD.

lancesix · 2026-05-07T13:57:43Z

+constexpr int MAX_GPUS = 8;
+constexpr int N_PER_DEVICE = 256;
+constexpr unsigned int group_size = 64;
+constexpr unsigned int num_groups = 2;


Why mix signed and unsigned?

lancesix · 2026-05-07T13:58:54Z

+/* Use small but non-trivial launch dimensions so the test runs quickly
+   while still creating multiple waves participating in the GWS
+   barrier.  GROUP_SIZE = 64 means 1 wave on wave64 architectures and
+   2 waves on wave32 architectures, for a total of 2 (or 4) waves.


can,t we check at runtime if we have wave32 or wave64? Could also build the testsuite with -mwavefrontsize64 to force wave64.

lancesix · 2026-05-07T15:12:48Z

+if {[build_executable "failed to prepare" $testfile $srcfile {debug hip}]} {
+    return
+}
+


We do not support debugging cooperative groups on gfx110x (see documented limitations). Testing this on gfx1100, I see:

FAIL: gdb.rocm/coop-group-grid-sync.exp: test_break_around_grid_sync: stop before grid.sync (the program exited) FAIL: gdb.rocm/coop-group-grid-sync.exp: test_threads_in_coop_kernel: stop in cooperative kernel (the program exited)

The testcase should probably detect for those devices and issue a XFAIL.

spatrang requested review from Copilot and lumachad May 7, 2026 13:08

Copilot started reviewing on behalf of spatrang May 7, 2026 13:10 View session

spatrang marked this pull request as ready for review May 7, 2026 13:14

spatrang requested a review from a team as a code owner May 7, 2026 13:14

Copilot AI reviewed May 7, 2026

View reviewed changes

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

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

Add dejagnu tests for cooperative group GWS debugging

33f1926

spatrang force-pushed the users/spatrang/coop-group-gws-tests branch from af78fac to 33f1926 Compare May 7, 2026 13:35

lancesix reviewed May 7, 2026

View reviewed changes

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Add dejagnu tests for cooperative group GWS debugging#116

Add dejagnu tests for cooperative group GWS debugging#116
spatrang wants to merge 1 commit intoamd-stagingfrom
users/spatrang/coop-group-gws-tests

spatrang commented May 7, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

Uh oh!

Uh oh!

lancesix left a comment

Uh oh!

lancesix May 7, 2026

Uh oh!

lancesix May 7, 2026

Uh oh!

lancesix May 7, 2026

Uh oh!

lancesix May 7, 2026

Uh oh!

lancesix May 7, 2026

Uh oh!

lancesix May 7, 2026

Uh oh!

lancesix May 7, 2026

Uh oh!

lancesix May 7, 2026

Uh oh!

lancesix May 7, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

spatrang commented May 7, 2026

Summary

Tests added

What gets verified

Skip / unsupported handling

Out of scope / follow-ups

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Uh oh!

Uh oh!

lancesix left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants