diff --git a/gdb/amd-dbgapi-target.c b/gdb/amd-dbgapi-target.c index f3bab7e60cc..4d752ab0e5c 100644 --- a/gdb/amd-dbgapi-target.c +++ b/gdb/amd-dbgapi-target.c @@ -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; @@ -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; } @@ -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); @@ -661,10 +672,10 @@ flatid_to_id (size_t flatid, const vec3_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> +partial_workgroup_sizes (thread_info *tp) { wave_info &info = get_thread_wave_info (tp); if (info.coords.dispatch_id == AMD_DBGAPI_DISPATCH_NONE) @@ -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 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 partial_wg_sizes; @@ -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. */ @@ -1183,6 +1211,21 @@ amd_dbgapi_target::grid_sizes (thread_info *thr) != AMD_DBGAPI_STATUS_SUCCESS) return std::nullopt; + std::array 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, @@ -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; } @@ -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) { @@ -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 work_group_sizes; if ((status = amd_dbgapi_dispatch_get_info ( @@ -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); @@ -4574,6 +4627,8 @@ info_dispatches_command (const char *args, int from_tty) ui_left, "target-id", "Target Id"); uiout->table_header (std::max (4, max_grid_width), ui_left, "grid", "Grid"); + uiout->table_header (std::max (7, max_cluster_width), + ui_left, "cluster", "Cluster"); uiout->table_header (std::max (9, max_workgroup_width), ui_left, "workgroup", "Workgroup"); uiout->table_header (7, ui_left, "fence", "Fence"); @@ -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 work_group_sizes; if ((status = amd_dbgapi_dispatch_get_info ( diff --git a/gdb/amd-dbgapi-target.h b/gdb/amd-dbgapi-target.h index b599eb7684f..543e3cbfa13 100644 --- a/gdb/amd-dbgapi-target.h +++ b/gdb/amd-dbgapi-target.h @@ -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> partial_workgroup_sizes + (thread_info *tp); + #endif /* GDB_AMD_DBGAPI_TARGET_H */ diff --git a/gdb/amdgpu-tdep.c b/gdb/amdgpu-tdep.c index cceee91dad7..98708950f18 100644 --- a/gdb/amdgpu-tdep.c +++ b/gdb/amdgpu-tdep.c @@ -1958,30 +1958,13 @@ 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); @@ -1989,20 +1972,8 @@ amdgpu_used_lanes_count (struct gdbarch *gdbarch, thread_info *tp) 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 (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); diff --git a/gdb/doc/gdb.texinfo b/gdb/doc/gdb.texinfo index 4e80a48571c..12858d3f5b4 100644 --- a/gdb/doc/gdb.texinfo +++ b/gdb/doc/gdb.texinfo @@ -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 @@ -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 @@ -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 @@ -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 @@ -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. @@ -30033,7 +30053,7 @@ 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 @@ -30041,11 +30061,19 @@ 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. diff --git a/gdb/testsuite/gdb.rocm/cluster.cpp b/gdb/testsuite/gdb.rocm/cluster.cpp new file mode 100644 index 00000000000..825b80acd59 --- /dev/null +++ b/gdb/testsuite/gdb.rocm/cluster.cpp @@ -0,0 +1,49 @@ +/* Copyright (C) 2026 Free Software Foundation, Inc. + Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved. + + This file is part of GDB. + + This program is free software; you can redistribute it and/or modify + it under the terms of the GNU General Public License as published by + the Free Software Foundation; either version 3 of the License, or + (at your option) any later version. + + This program is distributed in the hope that it will be useful, + but WITHOUT ANY WARRANTY; without even the implied warranty of + MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + GNU General Public License for more details. + + You should have received a copy of the GNU General Public License + along with this program. If not, see . */ + +/* Run a HIP program where workgroups are organized in clusters. */ + +#include +#include +#include +#include + +#include "rocm-test-utils.h" + +__global__ void +#ifdef WITH_CLUSTER +__cluster_dims__(3, 2, 1) +#endif +kernel () +{ + namespace cg = cooperative_groups; + cg::cluster_group c = cg::this_cluster(); + + /* Index of the calling block within cluster. */ + auto block_in_cluster = c.block_index (); + + __builtin_amdgcn_s_sleep (1); /* break-here */ +} + +int +main () +{ + kernel<<>>(); + CHECK (hipDeviceSynchronize ()); + return 0; +} diff --git a/gdb/testsuite/gdb.rocm/cluster.exp b/gdb/testsuite/gdb.rocm/cluster.exp new file mode 100644 index 00000000000..114a281294a --- /dev/null +++ b/gdb/testsuite/gdb.rocm/cluster.exp @@ -0,0 +1,107 @@ +# Copyright 2026 Free Software Foundation, Inc. +# Copyright (C) 2026 Advanced Micro Devices, Inc. All rights reserved. + +# This file is part of GDB. + +# This program is free software; you can redistribute it and/or modify +# it under the terms of the GNU General Public License as published by +# the Free Software Foundation; either version 3 of the License, or +# (at your option) any later version. + +# This program is distributed in the hope that it will be useful, +# but WITHOUT ANY WARRANTY; without even the implied warranty of +# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the +# GNU General Public License for more details. + +# You should have received a copy of the GNU General Public License +# along with this program. If not, see . + +# Test showing cluster coordinates. + +load_lib rocm.exp + +require allow_hip_tests + +standard_testfile .cpp + +if {[hip_device_is_less_than gfx1250]} { + unsupported "not testing clusters on device older than gfx1250" + return +} + +# First, build without clusters and check the dimensions in "info +# dispatch". + +if {[build_executable "failed to prepare" $testfile $srcfile {debug hip}]} { + return +} + +clean_restart $testfile + +with_rocm_gpu_lock { + with_test_prefix "without-cluster" { + if {![runto_main]} { + return + } + + set bp_line [gdb_get_line_number "break-here"] + gdb_breakpoint "$bp_line" -allow-pending -temporary + gdb_continue_to_breakpoint "break-here" + + # Expect grid and workgroup dimensions. + # Cluster dimensions is "-". + gdb_test "info dispatches" \ + "\\\[270,20,2\\\].*\-.*\\\[30,5,1\\\].*" + } +} + +# Now build with clusters, check the coordinates and dimensions. + +if {[build_executable "failed to prepare" $testfile $srcfile \ + {debug hip additional_flags=-DWITH_CLUSTER=1}]} { + return +} + +clean_restart $testfile + +with_rocm_gpu_lock { + if {![runto_main]} { + return + } + + set bp_line [gdb_get_line_number "break-here"] + gdb_breakpoint "$bp_line" -allow-pending -temporary + gdb_continue_to_breakpoint "break-here" + + set dispatch_pos "$decimal:$decimal:$decimal:$decimal" + set triple "\\(($decimal),($decimal),($decimal)\\)" + foreach_with_prefix wave [info_thread_get_wave_list] { + gdb_test "thread $wave" "Switching.*" "switch wave" + + gdb_test_multiple "info thread $wave" "capture coord" { + -re -wrap "AMDGPU Wave $dispatch_pos $triple/$triple/$decimal.*" { + set cluster_x $expect_out(1,string) + set cluster_y $expect_out(2,string) + set cluster_z $expect_out(3,string) + set wg_x $expect_out(4,string) + set wg_y $expect_out(5,string) + set wg_z $expect_out(6,string) + verbose -log "cluster: {$cluster_x,$cluster_y,$cluster_z}" + verbose -log "workgroup: {$wg_x,$wg_y,$wg_z}" + + gdb_test "print \$ttmp9" "= $cluster_x" "ttmp9 is cluster x" + gdb_test "print \$ttmp7 & 0xFFFF" "= $cluster_y" \ + "ttmp7 lo is cluster y" + gdb_test "print \$ttmp7 >> 16" "= $cluster_z" \ + "ttmp7 lo is cluster z" + + gdb_test "print block_in_cluster" \ + "= \{x = $wg_x, y = $wg_y, z = $wg_z\}" + } + } + } + + # Expect grid, cluster, and workgroup dimensions. + gdb_test "info dispatches" \ + "\\\[270,20,2\\\].*\\\[90,10,1\\\].*\\\[30,5,1\\\].*" +} diff --git a/gdb/testsuite/gdb.rocm/info-dispatches.exp b/gdb/testsuite/gdb.rocm/info-dispatches.exp index e9f0f2cc5e0..be17cde98d5 100644 --- a/gdb/testsuite/gdb.rocm/info-dispatches.exp +++ b/gdb/testsuite/gdb.rocm/info-dispatches.exp @@ -95,13 +95,13 @@ with_rocm_gpu_lock { set seen_header 0 set leading_star 0 set dispatch_ids [list] - set table_header_regex "\[ \t\]+Target\[ \t\]+Id\[ \t\]+Grid\[ \t\]+Workgroup\[ \t\]+Fence\[ \t\]+Kernel\[ \t\]+Function" + set table_header_regex "\[ \t\]+Target\[ \t\]+Id\[ \t\]+Grid\[ \t\]+Cluster\[ \t\]+Workgroup\[ \t\]+Fence\[ \t\]+Kernel\[ \t\]+Function" # Sample output: # - # Id Target Id Grid Workgroup Fence Kernel Function - # * 1 AMDGPU Dispatch 1:1:1 (PKID 1) [1,1,1] [1,1,1] B|Aa|Ra single_wave_kernel1() - # 2 AMDGPU Dispatch 1:3:2 (PKID 0) [1,1,1] [1,1,1] B|Aa|Ra single_wave_kernel2() + # Id Target Id Grid Cluster Workgroup Fence Kernel Function + # * 1 AMDGPU Dispatch 1:1:1 (PKID 1) [1,1,1] - [1,1,1] B|Aa|Ra single_wave_kernel1() + # 2 AMDGPU Dispatch 1:3:2 (PKID 0) [1,1,1] - [1,1,1] B|Aa|Ra single_wave_kernel2() # gdb_test_multiple "info dispatches" "info dispatches with multiple dispatches" -lbl { -re "$table_header_regex" { @@ -130,15 +130,15 @@ with_rocm_gpu_lock { set seen_header 0 set leading_star 0 set dispatch_ids [list] - set table_header_regex "\[ \t\]+Target\[ \t\]+Id\[ \t\]+Grid\[ \t\]+Workgroup\[ \t\]+Fence\[ \t\]+Address\[ \t\]+Spaces\[ \t\]+Kernel\[ \t\]+Descriptor\[ \t\]+Kernel\[ \t\]+Args\[ \t\]+Completion\[ \t\]+Signal\[ \t\]+Kernel\[ \t\]+Function" + set table_header_regex "\[ \t\]+Target\[ \t\]+Id\[ \t\]+Grid\[ \t\]+Cluster\[ \t\]+Workgroup\[ \t\]+Fence\[ \t\]+Address\[ \t\]+Spaces\[ \t\]+Kernel\[ \t\]+Descriptor\[ \t\]+Kernel\[ \t\]+Args\[ \t\]+Completion\[ \t\]+Signal\[ \t\]+Kernel\[ \t\]+Function" # Save kernel1 dispatch id for the next test. set kernel1_dispatch 0 # Sample output: - # Id Target Id Grid Workgroup Fence Address Spaces Kernel Descriptor Kernel Args Completion Signal Kernel Function - # * 1 AMDGPU Dispatch 1:1:1 (PKID 0) [1,1,1] [1,1,1] B|Aa|Ra Shared(0), Private(0) 0x00007ffff5dd8740 0x00007ffee0c00000 (nil) single_wave_kernel1() - # 2 AMDGPU Dispatch 1:3:2 (PKID 0) [1,1,1] [1,1,1] B|Aa|Ra Shared(0), Private(0) 0x00007ffff5dd8780 0x00007ffee0200000 (nil) single_wave_kernel2() + # 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) [1,1,1] - [1,1,1] B|Aa|Ra Shared(0), Private(0) 0x00007ffff5dd8740 0x00007ffee0c00000 (nil) single_wave_kernel1() + # 2 AMDGPU Dispatch 1:3:2 (PKID 0) [1,1,1] - [1,1,1] B|Aa|Ra Shared(0), Private(0) 0x00007ffff5dd8780 0x00007ffee0200000 (nil) single_wave_kernel2() # gdb_test_multiple "info dispatches -full" "info dispatches with full arg" -lbl { -re "$table_header_regex" { @@ -168,12 +168,12 @@ with_rocm_gpu_lock { # KERNEL1_DISPATCH. set seen_header 0 - set table_header_regex "\[ \t\]+Target\[ \t\]+Id\[ \t\]+Grid\[ \t\]+Workgroup\[ \t\]+Fence\[ \t\]+Kernel\[ \t\]+Function" + set table_header_regex "\[ \t\]+Target\[ \t\]+Id\[ \t\]+Grid\[ \t\]+Cluster\[ \t\]+Workgroup\[ \t\]+Fence\[ \t\]+Kernel\[ \t\]+Function" # Sample output: # - # Id Target Id Grid Workgroup Fence Kernel Function - # * 1 AMDGPU Dispatch 1:1:1 (PKID 1) [1,1,1] [1,1,1] B|Aa|Ra single_wave_kernel1() + # Id Target Id Grid Cluster Workgroup Fence Kernel Function + # * 1 AMDGPU Dispatch 1:1:1 (PKID 1) [1,1,1] - [1,1,1] B|Aa|Ra single_wave_kernel1() # gdb_test_multiple "info dispatches $kernel1_dispatch" \ "info dispatches given a dispatch id" -lbl {