handle clusters by aktemur · Pull Request #206 · ROCm/ROCgdb

aktemur · 2026-07-14T12:39:23Z

Motivation

New feature: workgroup clusters

Technical Details

Workgroups can be organized in clusters. dbgapi provides new query functions. Use these queries to show cluster information.

Test Plan

A basic test is included.

aktemur · 2026-07-16T15:41:03Z

Here is a sample output of "info threads" with clusters:

info threads
  Id   Target Id                                   Frame 
...
* 6    AMDGPU Wave 1:2:1:1 (0,0,0)/(0,0,0)/0       kernel () at ...
  7    AMDGPU Wave 1:2:1:2 (0,0,0)/(0,0,0)/1       kernel () at ...
...
  36   AMDGPU Wave 1:2:1:31 (2,0,1)/(0,0,0)/0      kernel () at ...
  37   AMDGPU Wave 1:2:1:32 (2,0,1)/(0,0,0)/1      kernel () at ...
...
  71   AMDGPU Wave 1:2:1:66 (1,0,0)/(1,0,0)/0      kernel () at ...
...
  95   AMDGPU Wave 1:2:1:90 (1,0,0)/(2,1,0)/4      kernel () at ...

And "info dispatches":

info dispatches
  Id   Target Id                      Grid       Cluster   Workgroup Fence   Kernel Function 
* 1    AMDGPU Dispatch 1:2:1 (PKID 1) [270,20,2] [90,10,1] [30,5,1]  B|As|Rs 0x00007ffff6084000

The function 'amdgpu_used_lanes_count' in amdgpu-tdep.c and the function 'lane_workgroup_pos' in amd-dbgapi-target.c both compute partial workgroup sizes the same way. Extract the common code into a function named 'partial_workgroup_sizes' and reuse. This is a refactoring with no behavioral change.

aktemur · 2026-07-28T07:38:56Z

Test is updated. It now includes synchronization and checks blockIdx dimensions as obtained from the program to those calculated based on cluster&wg coordinates.

lancesix · 2026-07-28T09:07:50Z


 @smallexample
-AMDGPU Wave @var{agent-id}:@var{queue-id}:@var{dispatch-id}:@var{wave-id} (@var{work-group-x},@var{work-group-y},@var{work-group-z})/@var{work-group-thread-index}
+AMDGPU Wave @var{agent-id}:@var{queue-id}:@var{dispatch-id}:@var{wave-id} (@var{cluster-x},@var{cluster-y},@var{cluster-z})/(@var{work-group-x},@var{work-group-y},@var{work-group-z})/@var{work-group-thread-index}


just random thought.

When imagining this initially, I kind of assumed we would use something like {cx,cy,cz}(x,y,z)/n with cluster dimensions in another set of brackets ({}). Not to say I mind the cluster/wg/wave approach, I am not sure I have a strong opinion at this point, but wanted to drop this as a note.

just random thought.

When imagining this initially, I kind of assumed we would use something like {cx,cy,cz}(x,y,z)/n with cluster dimensions in another set of brackets ({}). Not to say I mind the cluster/wg/wave approach, I am not sure I have a strong opinion at this point, but wanted to drop this as a note.

I have no objections to using braces.

lancesix · 2026-07-28T09:11:33Z

+  int bz = blockIdx.z;
+
+  /* Define a barrier to also test debuggability in the existence of
+     synchronization.  */


the interesting debugability part here would be if some waves / workgroups reached the barrier when some other have not. Seems that in this case, breaking everyone after the barrier limits the probabilities of having some waves before and some after the barrier.

the interesting debugability part here would be if some waves / workgroups reached the barrier when some other have not. Seems that in this case, breaking everyone after the barrier limits the probabilities of having some waves before and some after the barrier.

I don't fully follow. In principle, when we break at or before the barrier, there is always a possibility that some waves are behind. Only after the barrier we have the guarantee that all waves are now beyond the barrier. That's the purpose of having a barrier, no?
I mean, if I break before the barrier and not all waves are there, why should that be a bug? It would be just a timing issue.

The isa has a split barrier model, where waves need to notify when they arrive at the barrier, then wait on it. The barrier has an internal state counting how many waves have signaled (s_barrier_signal), and will keep signaled waves from continuing (s_barrier_wait) until all have signaled. During debug events, we trigger context save. The main thing (which has had issues in the past) which could get wrong during context save is not saving / restoring the barrier state correctly. This is what is tested by gdb.rocm/device-barrier.exp at the workgroup level. I think my point is that while at introducing cluster concepts, it might be worth introducing a similar test exercising cluster barrier save and restore.

Ok, thanks for the clarification. The cluster API has barrier_arrive and barrier_wait separately, for which there is a comment for sync:

// Sync the cluster, equivalent to c.barrier_wait(c.barrier_arrive());

How about having barrier_arrive and barrier_wait, with a breakpoint in-between? That sounds like it should satisfy the previously problematic scenario.

A good test would probably test the following cases:

some of the waves have called arrived, and are now waiting (or about to)

all the waves have called arrived, but none have called wait

all have called arrived, some have called wait (so might be past the barrier)

but maybe it feels this particular matrix would be in a different test as checking the workgroup cluster dimensions.

but maybe it feels this particular matrix would be in a different test as checking the workgroup cluster dimensions.

Agree.

lancesix

also, shouldn't we have some configury stuff to check dbgapi's version to be sure AMD_DBGAPI_DISPATCH_INFO_CLUSTER_SIZES is provided?

aktemur · 2026-07-28T11:35:12Z

also, shouldn't we have some configury stuff to check dbgapi's version to be sure AMD_DBGAPI_DISPATCH_INFO_CLUSTER_SIZES is provided?

In nowhere else we have a similar check. I thought we'd first merge dbgapi PR, which bumps the version, and then build ROCgdb with that version, which brings us the guard here:

  if (major != AMD_DBGAPI_VERSION_MAJOR || minor < AMD_DBGAPI_VERSION_MINOR)
    error (_("amd-dbgapi library version mismatch, got %d.%d.%d, need %d.%d+"),
           major, minor, patch, AMD_DBGAPI_VERSION_MAJOR,
           AMD_DBGAPI_VERSION_MINOR);

lancesix · 2026-07-28T11:43:17Z

also, shouldn't we have some configury stuff to check dbgapi's version to be sure AMD_DBGAPI_DISPATCH_INFO_CLUSTER_SIZES is provided?

In nowhere else we have a similar check. I thought we'd first merge dbgapi PR, which bumps the version, and then build ROCgdb with that version, which brings us the guard here:
  if (major != AMD_DBGAPI_VERSION_MAJOR || minor < AMD_DBGAPI_VERSION_MINOR)
    error (_("amd-dbgapi library version mismatch, got %d.%d.%d, need %d.%d+"),
           major, minor, patch, AMD_DBGAPI_VERSION_MAJOR,
           AMD_DBGAPI_VERSION_MINOR);

We have https://github.com/ROCm/ROCgdb/blob/amd-staging/gdb/configure.ac#L330-L331 which currently guards building GDB with 0.80 or more.

Either we wait for dbgapi to land, and when we have the right version bump configure.ac in this PR, or have something more dynamic (we do not have this yet) where configure can test dbgapi for this feature (use AC_COMPILE_IFELSE to try to build a program including dbgapi and using this new API, and set a compile flag based on that). I don't think I have a strong opinion either way, the AC_COMPILE_IFELSE could allow merging GDB first.

aktemur · 2026-07-28T11:48:52Z

Either we wait for dbgapi to land, and when we have the right version bump configure.ac in this PR, or have something more dynamic (we do not have this yet) where configure can test dbgapi for this feature (use AC_COMPILE_IFELSE to try to build a program including dbgapi and using this new API, and set a compile flag based on that). I don't think I have a strong opinion either way, the AC_COMPILE_IFELSE could allow merging GDB first.

Ok, gotcha. I'll go with incrementing the configure version in this PR and leave the dynamic detection to a separate feature task.

aktemur · 2026-07-28T12:40:12Z

The last push updates the document, removes c.sync () from the test (will be handled separately), bumps the version check in configure.

Starting with gfx1250, workgroups may be organized into clusters, giving a workitem -> wave -> workgroup -> cluster -> grid hierarchy. Handle clusters by taking them into account when calculating workitem/lane coordinates, when showing wave target ids, and in the "info dispatches" command.

aktemur · 2026-07-31T10:29:47Z

The last push converts in-grid group ids to in-cluster group ids.

aktemur requested a review from a team as a code owner July 14, 2026 12:39

github-actions Bot assigned lumachad Jul 14, 2026

aktemur assigned lancesix and unassigned lumachad Jul 14, 2026

aktemur force-pushed the users/aktemur/cluster branch 2 times, most recently from 5b5bb47 to 5af76b9 Compare July 16, 2026 07:55

aktemur mentioned this pull request Jul 16, 2026

feat(rocdbgapi): introduce clusters of workgroups ROCm/rocm-systems#7992

Open

aktemur force-pushed the users/aktemur/cluster branch from 5af76b9 to 4c9d135 Compare July 28, 2026 07:37

lancesix reviewed Jul 28, 2026

View reviewed changes

aktemur force-pushed the users/aktemur/cluster branch from 4c9d135 to 2d1a559 Compare July 28, 2026 12:39

aktemur force-pushed the users/aktemur/cluster branch from 2d1a559 to 1c0dea6 Compare July 31, 2026 10:28

Uh oh!

Conversation

aktemur commented Jul 14, 2026

Motivation

Technical Details

Test Plan

Uh oh!

aktemur commented Jul 16, 2026

Uh oh!

aktemur commented Jul 28, 2026

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

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!

lancesix left a comment

Choose a reason for hiding this comment

Uh oh!

aktemur commented Jul 28, 2026

Uh oh!

lancesix commented Jul 28, 2026

Uh oh!

aktemur commented Jul 28, 2026

Uh oh!

aktemur commented Jul 28, 2026

Uh oh!

aktemur commented Jul 31, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants