Skip to content

handle clusters - #206

Open
aktemur wants to merge 2 commits into
amd-stagingfrom
users/aktemur/cluster
Open

handle clusters#206
aktemur wants to merge 2 commits into
amd-stagingfrom
users/aktemur/cluster

Conversation

@aktemur

@aktemur aktemur commented Jul 14, 2026

Copy link
Copy Markdown
Contributor

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
aktemur requested a review from a team as a code owner July 14, 2026 12:39
@aktemur aktemur assigned lancesix and unassigned lumachad Jul 14, 2026
@aktemur
aktemur force-pushed the users/aktemur/cluster branch 2 times, most recently from 5b5bb47 to 5af76b9 Compare July 16, 2026 07:55
@aktemur aktemur mentioned this pull request Jul 16, 2026
Open
@aktemur

aktemur commented Jul 16, 2026

Copy link
Copy Markdown
Contributor Author

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

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

aktemur commented Jul 28, 2026

Copy link
Copy Markdown
Contributor Author

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
lancesix reviewed Jul 28, 2026
View reviewed changes
Comment thread gdb/doc/gdb.texinfo Outdated
Comment thread gdb/doc/gdb.texinfo
Comment thread gdb/testsuite/gdb.rocm/cluster.cpp Outdated
lancesix
lancesix reviewed Jul 28, 2026
View reviewed changes

@lancesix lancesix left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

aktemur commented Jul 28, 2026

Copy link
Copy Markdown
Contributor Author

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

lancesix commented Jul 28, 2026

Copy link
Copy Markdown
Collaborator

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

aktemur commented Jul 28, 2026

Copy link
Copy Markdown
Contributor Author

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
aktemur force-pushed the users/aktemur/cluster branch from 4c9d135 to 2d1a559 Compare July 28, 2026 12:39
@aktemur

aktemur commented Jul 28, 2026

Copy link
Copy Markdown
Contributor Author

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

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

aktemur commented Jul 31, 2026

Copy link
Copy Markdown
Contributor Author

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

@aktemur
gdb, amdgpu: extract function 'partial_workgroup_sizes'
fe8e90f 
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
aktemur force-pushed the users/aktemur/cluster branch from 1c0dea6 to 6abe945 Compare August 12, 2026 12:36
@aktemur

aktemur commented Aug 12, 2026

Copy link
Copy Markdown
Contributor Author

The force-push above is a rebase.

@aktemur
aktemur force-pushed the users/aktemur/cluster branch from 6abe945 to 5afc339 Compare August 12, 2026 14:57
@aktemur

aktemur commented Aug 12, 2026

Copy link
Copy Markdown
Contributor Author

The force-push above uses AMD_DBGAPI_WAVE_INFO_WORKGROUP_COORD_IN_CLUSTER to get the wg coordinate within the cluster.

@aktemur
gdb, amdgpu: support workgroup clusters
b56c366 
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
aktemur force-pushed the users/aktemur/cluster branch from 5afc339 to b56c366 Compare August 13, 2026 11:27
@aktemur

aktemur commented Aug 13, 2026

Copy link
Copy Markdown
Contributor Author

The update above adds a new test that uses barrier arrive/wait.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@lancesix lancesix lancesix left review comments

At least 1 approving review is required to merge this pull request.

Assignees

@lancesix lancesix

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

3 participants

@aktemur @lancesix @lumachad