Skip to content

handle clusters (alternative) - #287

Merged
aktemur merged 2 commits into
amd-stagingfrom
users/aktemur/cluster-111
Oct 6, 2026
Merged

aktemur merged 2 commits into
amd-stagingfrom
users/aktemur/cluster-111

Conversation

@aktemur

@aktemur aktemur commented Aug 18, 2026

Copy link
Copy Markdown
Contributor

This is the same as PR #206, except that it uses the alternative dbgapi in PR ROCm/rocm-systems#10241.

@aktemur
aktemur requested a review from a team as a code owner August 18, 2026 09:16
@lumachad lumachad removed their assignment Aug 19, 2026
@aktemur
aktemur force-pushed the users/aktemur/cluster-111 branch from 7f5823a to cc43daa Compare September 22, 2026 08:47
@aktemur

aktemur commented Sep 22, 2026

Copy link
Copy Markdown
Contributor Author

The force-push above is a rebase.

@aktemur
aktemur force-pushed the users/aktemur/cluster-111 branch 2 times, most recently from 51fbc85 to d08a951 Compare September 28, 2026 07:58
Comment thread gdb/amd-dbgapi-target.h Resolved
Comment thread gdb/doc/gdb.texinfo Resolved
Comment thread gdb/doc/gdb.texinfo Resolved
Comment thread gdb/testsuite/gdb.rocm/cluster.exp Resolved
Comment thread gdb/amd-dbgapi-target.c Resolved
Comment thread gdb/amd-dbgapi-target.c Resolved
Comment thread gdb/amd-dbgapi-target.c Resolved
Comment thread gdb/amd-dbgapi-target.c Resolved
@lancesix lancesix assigned aktemur and unassigned lancesix Sep 28, 2026
@aktemur
aktemur force-pushed the users/aktemur/cluster-111 branch from d08a951 to 35a15f6 Compare September 29, 2026 15:45
@aktemur

aktemur commented Sep 29, 2026

Copy link
Copy Markdown
Contributor Author

The force-push above is a rebase.

@aktemur
aktemur force-pushed the users/aktemur/cluster-111 branch from 35a15f6 to a78e1c9 Compare September 29, 2026 16:35
@aktemur aktemur assigned lancesix and unassigned aktemur Sep 29, 2026
@aktemur
aktemur force-pushed the users/aktemur/cluster-111 branch from a78e1c9 to 621e132 Compare October 5, 2026 07:44
@aktemur

aktemur commented Oct 6, 2026

Copy link
Copy Markdown
Contributor Author

CHANGELOG needs to be updated, too.

@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.

just one minor thing which could end-up annoying when testing with -O1 and above, otherwise seems reasonable to me.

Comment thread gdb/testsuite/gdb.rocm/cluster.cpp Resolved
@aktemur
aktemur force-pushed the users/aktemur/cluster-111 branch from 621e132 to 2aad83f Compare October 6, 2026 13:23
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.
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.

Bug: AIROCGDB-77
@aktemur
aktemur force-pushed the users/aktemur/cluster-111 branch from 2aad83f to a6bc758 Compare October 6, 2026 14:58
@aktemur
aktemur merged commit a6bc758 into amd-staging Oct 6, 2026
36 checks passed
@aktemur
aktemur deleted the users/aktemur/cluster-111 branch October 6, 2026 16:00
@aktemur aktemur mentioned this pull request Oct 8, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants