Skip to content

handle clusters #206

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
aktemur wants to merge 2 commits into
base: amd-staging
Choose a base branch
from
+295 −68
Open

handle clusters #206

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
21c65e7
gdb, amdgpu: extract function 'partial_workgroup_sizes'
aktemur Jul 15, 2026
5af76b9
gdb, amdgpu: support workgroup clusters
aktemur Jul 15, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
92 changes: 79 additions & 13 deletions gdb/amd-dbgapi-target.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -157,6 +157,7 @@ struct wave_coordinates
amd_dbgapi_dispatch_id_t dispatch_id = AMD_DBGAPI_DISPATCH_NONE;
amd_dbgapi_queue_id_t queue_id = AMD_DBGAPI_QUEUE_NONE;
amd_dbgapi_agent_id_t agent_id = AMD_DBGAPI_AGENT_NONE;
vec3_u32_t cluster_ids {UINT32_MAX, UINT32_MAX, UINT32_MAX};
vec3_u32_t group_ids {UINT32_MAX, UINT32_MAX, UINT32_MAX};
uint32_t wave_in_group = UINT32_MAX;

Expand Down Expand Up @@ -476,6 +477,12 @@ wave_coordinates::workgroup_coord_str () const
pulongest (group_ids[1]), pulongest (group_ids[2]))
: "(?,?,?)");

if (cluster_ids[0] != UINT32_MAX)
str = string_printf ("(%s,%s,%s)/%s", pulongest (cluster_ids[0]),
pulongest (cluster_ids[1]),
pulongest (cluster_ids[2]),
str.c_str ());

return str;
}

Expand Down Expand Up @@ -520,6 +527,10 @@ wave_coordinates::fetch ()
AMD_DBGAPI_WAVE_INFO_WORKGROUP_COORD,
sizeof (group_ids), &group_ids);

amd_dbgapi_wave_get_info (wave_id,
AMD_DBGAPI_WAVE_INFO_CLUSTER_COORD,
sizeof (cluster_ids), &cluster_ids);

amd_dbgapi_wave_get_info (wave_id,
AMD_DBGAPI_WAVE_INFO_WAVE_NUMBER_IN_WORKGROUP,
sizeof (wave_in_group), &wave_in_group);
Expand Down Expand Up @@ -661,10 +672,10 @@ flatid_to_id (size_t flatid, const vec3_t<size_t> &sizes)
return coord_id;
}

/* Return the work-group position of the work-item assigned to lane LANE. */
/* See amd-dbgapi-target.h. */

static opt_vec3_u32_t
lane_workgroup_pos (thread_info *tp, int lane)
std::optional<std::array<size_t, 3>>
partial_workgroup_sizes (thread_info *tp)
{
wave_info &info = get_thread_wave_info (tp);
if (info.coords.dispatch_id == AMD_DBGAPI_DISPATCH_NONE)
Expand All @@ -676,18 +687,19 @@ lane_workgroup_pos (thread_info *tp, int lane)
}

vec3_u32_t grid_sizes;
dispatch_get_info_throw (info.coords.dispatch_id,
AMD_DBGAPI_DISPATCH_INFO_GRID_SIZES,
grid_sizes);
/* If in cluster mode, use cluster sizes instead of the grid,
because the cluster in this case becomes the workgroup
container. */
auto query = ((info.coords.cluster_ids[0] == UINT32_MAX)
? AMD_DBGAPI_DISPATCH_INFO_GRID_SIZES
: AMD_DBGAPI_DISPATCH_INFO_CLUSTER_SIZES);
dispatch_get_info_throw (info.coords.dispatch_id, query, grid_sizes);

vec3_t<uint16_t> work_group_sizes;
dispatch_get_info_throw (info.coords.dispatch_id,
AMD_DBGAPI_DISPATCH_INFO_WORKGROUP_SIZES,
work_group_sizes);

size_t lane_count;
wave_get_info_throw (tp, AMD_DBGAPI_WAVE_INFO_LANE_COUNT, lane_count);

/* Find the work-group item sizes for each axis, taking into account
the work-items that actually fit in the grid. */
vec3_t<size_t> partial_wg_sizes;
Expand All @@ -701,9 +713,25 @@ lane_workgroup_pos (thread_info *tp, int lane)
partial_wg_sizes[i] = work_item_end - work_item_start;
}

return partial_wg_sizes;
}

/* Return the work-group position of the work-item assigned to lane LANE. */

static opt_vec3_u32_t
lane_workgroup_pos (thread_info *tp, int lane)
{
auto partial_wg_sizes = partial_workgroup_sizes (tp);
if (!partial_wg_sizes.has_value ())
return std::nullopt;

size_t lane_count;
wave_get_info_throw (tp, AMD_DBGAPI_WAVE_INFO_LANE_COUNT, lane_count);

wave_info &info = get_thread_wave_info (tp);
size_t flatid = info.coords.wave_in_group * lane_count + lane;

return flatid_to_id (flatid, partial_wg_sizes);
return flatid_to_id (flatid, partial_wg_sizes.value ());
}

/* Return the lane's work-group position as a string. */
Expand Down Expand Up @@ -1183,6 +1211,21 @@ amd_dbgapi_target::grid_sizes (thread_info *thr)
!= AMD_DBGAPI_STATUS_SUCCESS)
return std::nullopt;

std::array<uint32_t, 3> cluster_sizes;
if (info.coords.cluster_ids[0] == UINT32_MAX
|| (amd_dbgapi_dispatch_get_info (info.coords.dispatch_id,
AMD_DBGAPI_DISPATCH_INFO_CLUSTER_SIZES,
sizeof (cluster_sizes),
cluster_sizes.data ())
!= AMD_DBGAPI_STATUS_SUCCESS))
{
/* If cluster info is not available, in the 3D hierarchy, the
"unit" under the grid is workgroup. */
cluster_sizes[0] = group_sizes[0];
cluster_sizes[1] = group_sizes[1];
cluster_sizes[2] = group_sizes[2];
}

vec3_u32_t grid_sizes;
if (amd_dbgapi_dispatch_get_info (info.coords.dispatch_id,
AMD_DBGAPI_DISPATCH_INFO_GRID_SIZES,
Expand All @@ -1193,7 +1236,7 @@ amd_dbgapi_target::grid_sizes (thread_info *thr)

/* Convert GRID_SIZES from "work-item" unit to "work-group" unit. */
for (size_t i = 0; i < 3; ++i)
grid_sizes[i] /= group_sizes[i];
grid_sizes[i] /= cluster_sizes[i];

return grid_sizes;
}
Expand Down Expand Up @@ -4478,7 +4521,7 @@ info_dispatches_command (const char *args, int from_tty)
{
size_t n_dispatches{ 0 }, max_target_id_width{ 0 },
max_grid_width{ 0 }, max_workgroup_width{ 0 },
max_address_spaces_width{ 0 };
max_cluster_width { 0 }, max_address_spaces_width{ 0 };

for (auto &&value : all_filtered_dispatches)
{
Expand Down Expand Up @@ -4512,6 +4555,16 @@ info_dispatches_command (const char *args, int from_tty)
= std::max (max_grid_width,
ndim_string (dims, grid_sizes).size ());

/* Clusters are optional, depending on architecture support. */
vec3_u32_t cluster_sizes;
if ((status = amd_dbgapi_dispatch_get_info (
dispatch_id, AMD_DBGAPI_DISPATCH_INFO_CLUSTER_SIZES,
sizeof (cluster_sizes), &cluster_sizes[0]))
== AMD_DBGAPI_STATUS_SUCCESS)
max_cluster_width
= std::max (max_cluster_width,
ndim_string (dims, cluster_sizes).size ());

/* workgroup */
vec3_t<uint16_t> work_group_sizes;
if ((status = amd_dbgapi_dispatch_get_info (
Expand Down Expand Up @@ -4563,7 +4616,7 @@ info_dispatches_command (const char *args, int from_tty)
}

/* Header: */
table_emitter.emplace (uiout, opts.full ? 11 : 7, n_dispatches,
table_emitter.emplace (uiout, opts.full ? 12 : 8, n_dispatches,
"InfoRocmDispatchesTable");
size_t addr_width = 2 + (gdbarch_ptr_bit (gdbarch) / 4);

Expand All @@ -4574,6 +4627,8 @@ info_dispatches_command (const char *args, int from_tty)
ui_left, "target-id", "Target Id");
uiout->table_header (std::max<size_t> (4, max_grid_width), ui_left,
"grid", "Grid");
uiout->table_header (std::max<size_t> (7, max_cluster_width),
ui_left, "cluster", "Cluster");
uiout->table_header (std::max<size_t> (9, max_workgroup_width),
ui_left, "workgroup", "Workgroup");
uiout->table_header (7, ui_left, "fence", "Fence");
Expand Down Expand Up @@ -4654,6 +4709,17 @@ info_dispatches_command (const char *args, int from_tty)

uiout->field_string ("grid", ndim_string (dims, grid_sizes));

/* cluster */
vec3_u32_t cluster_sizes;
if ((status = amd_dbgapi_dispatch_get_info (
dispatch_id, AMD_DBGAPI_DISPATCH_INFO_CLUSTER_SIZES,
sizeof (cluster_sizes), &cluster_sizes[0]))
!= AMD_DBGAPI_STATUS_SUCCESS)
uiout->field_string ("cluster", "-");
else
uiout->field_string ("cluster",
ndim_string (dims, cluster_sizes));

/* workgroup */
vec3_t<uint16_t> work_group_sizes;
if ((status = amd_dbgapi_dispatch_get_info (
Expand Down
6 changes: 6 additions & 0 deletions gdb/amd-dbgapi-target.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -163,4 +163,10 @@ dispatch_get_info_throw (amd_dbgapi_dispatch_id_t dispatch_id,
dispatch_id.handle, get_status_string (status));
}

/* Return the sizes of the workgroup that thread TP belongs to,
taking into account the work-items that actually fit in the grid. */

std::optional<std::array<size_t, 3>> partial_workgroup_sizes
(thread_info *tp);

#endif /* GDB_AMD_DBGAPI_TARGET_H */
39 changes: 5 additions & 34 deletions gdb/amdgpu-tdep.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1958,51 +1958,22 @@ amdgpu_supported_lanes_count (struct gdbarch *gdbarch, thread_info *tp)
static int
amdgpu_used_lanes_count (struct gdbarch *gdbarch, thread_info *tp)
{
amd_dbgapi_dispatch_id_t dispatch_id;
if (wave_get_info (tp, AMD_DBGAPI_WAVE_INFO_DISPATCH, dispatch_id)
!= AMD_DBGAPI_STATUS_SUCCESS)
auto wg_sizes = partial_workgroup_sizes (tp);
if (!wg_sizes.has_value ())
{
/* The dispatch associated with a wave is not available. A wave
may not have an associated dispatch if attaching to a process
with already existing waves. In that case, all we can do is
claim that all lanes are used. */
/* In this case, all we can do is claim that all lanes are used. */
return amdgpu_supported_lanes_count (gdbarch, tp);
}

uint32_t grid_sizes[3];
dispatch_get_info_throw (dispatch_id,
AMD_DBGAPI_DISPATCH_INFO_GRID_SIZES,
grid_sizes);

uint16_t work_group_sizes[3];
dispatch_get_info_throw (dispatch_id,
AMD_DBGAPI_DISPATCH_INFO_WORKGROUP_SIZES,
work_group_sizes);

uint32_t group_ids[3];
wave_get_info_throw (tp, AMD_DBGAPI_WAVE_INFO_WORKGROUP_COORD, group_ids);

uint32_t wave_in_group;
wave_get_info_throw (tp, AMD_DBGAPI_WAVE_INFO_WAVE_NUMBER_IN_WORKGROUP,
wave_in_group);

size_t lane_count;
wave_get_info_throw (tp, AMD_DBGAPI_WAVE_INFO_LANE_COUNT, lane_count);

size_t work_group_item_sizes[3];
for (int i = 0; i < 3; i++)
{
size_t item_start
= static_cast<size_t> (group_ids[i]) * work_group_sizes[i];
size_t item_end = item_start + work_group_sizes[i];
if (item_end > grid_sizes[i])
item_end = grid_sizes[i];
work_group_item_sizes[i] = item_end - item_start;
}

size_t work_items = (work_group_item_sizes[0]
* work_group_item_sizes[1]
* work_group_item_sizes[2]);
size_t work_items = (wg_sizes.value ()[0] * wg_sizes.value ()[1]
* wg_sizes.value ()[2]);

size_t work_items_left = work_items - wave_in_group * lane_count;
return std::min (work_items_left, lane_count);
Expand Down
48 changes: 38 additions & 10 deletions gdb/doc/gdb.texinfo
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24448,6 +24448,12 @@ The x, y, and z heterogeneous grid dimensions, in that order, in terms of
work-items of the heterogeneous dispatch. The number of dimensions displayed
matches the dimensionality of the heterogeneous dispatch.

@item
The x, y, and z heterogeneous cluster dimensions, in that order, in terms of
work-items of the heterogeneous dispatch. The number of dimensions displayed
matches the dimensionality of the heterogeneous dispatch. If the architecture
of the dispatch does not support clusters, a @code{-} is shown.

@item
The x, y, and z heterogeneous work-group dimensions, in that order, in terms of
work-items of the heterogeneous dispatch. The number of dimensions displayed
Expand Down Expand Up @@ -24494,8 +24500,8 @@ For example,

@smallexample
(@value{GDBP}) info dispatches -full
Id Target Id Grid Workgroup Fence Address Spaces Kernel Descriptor Kernel Args Completion Signal Kernel Function
* 1 AMDGPU Dispatch 1:1:1 (PKID 0) [256,1,1] [128,1,1] B|As Shared(0), Private(220) 0x00007ffde5409800 0x00007ffff7e00000 (nil) bit_extract_kernel(unsigned int*, unsigned int const*, unsigned long)
Id Target Id Grid Cluster Workgroup Fence Address Spaces Kernel Descriptor Kernel Args Completion Signal Kernel Function
* 1 AMDGPU Dispatch 1:1:1 (PKID 0) [256,1,1] - [128,1,1] B|As Shared(0), Private(220) 0x00007ffde5409800 0x00007ffff7e00000 (nil) bit_extract_kernel(unsigned int*, unsigned int const*, unsigned long)
@end smallexample

If you're debugging multiple inferiors, @value{GDBN} displays
Expand Down Expand Up @@ -29448,20 +29454,24 @@ The @acronym{AMD} @acronym{GPU} thread target identifier (@var{systag})
string has the following format:

@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}
@end smallexample

It is used in the @samp{Target ID} column of the @samp{info threads} command.
The cluster coordinates in the format are shown only if the work-group
is in a cluster.

@item Lane Target ID
The @acronym{AMD} @acronym{GPU} lane target identifier (@var{lane_systag})
string has the following format:

@smallexample
AMDGPU Lane @var{agent-id}:@var{queue-id}:@var{dispatch-id}:@var{wave-id}/@var{lane-index} (@var{work-group-x},@var{work-group-y},@var{work-group-z})[@var{work-item-x},@var{work-item-y},@var{work-item-z}]
AMDGPU Lane @var{agent-id}:@var{queue-id}:@var{dispatch-id}:@var{wave-id}/@var{lane-index} (@var{cluster-x},@var{cluster-y},@var{cluster-z})/(@var{work-group-x},@var{work-group-y},@var{work-group-z})[@var{work-item-x},@var{work-item-y},@var{work-item-z}]
@end smallexample

It is used in the @samp{Target ID} column of the @samp{info lanes} command.
The cluster coordinates in the format are shown only if the work-group
is in a cluster.

@end table

Expand All @@ -29472,10 +29482,11 @@ convenience variables:

@item $_dispatch_pos
The string returned by the @code{$_dispatch_pos} debugger convenience
variable has the following format:
variable has the following format, where cluster coordinates are included
only if the work-group is in a cluster.

@smallexample
(@var{work-group-x},@var{work-group-y},@var{work-group-z})/@var{work-group-thread-index}
(@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}
@end smallexample

@item $_thread_workgroup_pos
Expand Down Expand Up @@ -29518,10 +29529,19 @@ identifier identifier respectively. The identifiers are per process.
The @acronym{AMD} @acronym{GPU} target driver packet identifier. The
identifier is per queue.

@item cluster-x
@itemx cluster-y
@itemx cluster-z
The grid position of the thread's work-group's cluster within the
heterogeneous dispatch. This position is not shown if the work-group
does not belong to a cluster.

@item work-group-x
@itemx work-group-y
@itemx work-group-z
The grid position of the thread's work-group within the heterogeneous dispatch.
The position of the thread's work-group within the heterogeneous
dispatch, if the work-group belongs to a cluster. Otherwise
it is the position of the work-group within the grid.

@item work-group-thread-index
The thread's number within the heterogeneous work-group.
Expand Down Expand Up @@ -30033,19 +30053,27 @@ The debugger convenience variable @code{$_wave_id} is available which
returns a string that has the format:

@smallexample
(@var{work-group-x},@var{work-group-y},@var{work-group-z})/@var{work-group-thread-index}
(@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}
@end smallexample

@noindent
Where:

@table @var

@item cluster-x
@itemx cluster-y
@itemx cluster-z
The grid position of the thread's work-group's cluster within the
heterogeneous dispatch. This position is not shown if the work-group
does not belong to a cluster.

@item work-group-x
@itemx work-group-y
@itemx work-group-z
The grid position of the thread's work-group within the heterogeneous
dispatch.
The position of the thread's work-group within the heterogeneous
dispatch, if the work-group belongs to a cluster. Otherwise
it is the position of the work-group within the grid.

@item work-group-thread-index
The thread's number within the heterogeneous work-group.
Expand Down
Loading
Loading