handle clusters (alternative)

gdb/amd-dbgapi-target.c

@@ -157,8 +157,11 @@ 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 group_ids {UINT32_MAX, UINT32_MAX, UINT32_MAX};
+  vec3_u32_t cluster_dims {UINT32_MAX, UINT32_MAX, UINT32_MAX};
+  vec3_u32_t cluster_ids {UINT32_MAX, UINT32_MAX, UINT32_MAX};
+  vec3_u32_t group_ids_in_cluster {UINT32_MAX, UINT32_MAX, UINT32_MAX};
   uint32_t wave_in_group = UINT32_MAX;
+  bool in_cluster = false;
 
   explicit wave_coordinates (amd_dbgapi_wave_id_t wave_id)
     : wave_id (wave_id)
@@ -470,13 +473,28 @@ wave_coordinates::hierarchy_str () const
 std::string
 wave_coordinates::workgroup_coord_str () const
 {
-  std::string str
-    = (group_ids[0] != UINT32_MAX
-       ? string_printf ("(%s,%s,%s)", pulongest (group_ids[0]),
-			pulongest (group_ids[1]), pulongest (group_ids[2]))
+  std::string grid_pos
+    = (cluster_ids[0] != UINT32_MAX
+       ? string_printf ("(%s,%s,%s)", pulongest (cluster_ids[0]),
+			pulongest (cluster_ids[1]),
+			pulongest (cluster_ids[2]))
        : "(?,?,?)");
 
-  return str;
+  if (in_cluster)
+    {
+      std::string cluster_pos
+	= (group_ids_in_cluster[0] != UINT32_MAX
+	   ? string_printf ("(%s,%s,%s)",
+			    pulongest (group_ids_in_cluster[0]),
+			    pulongest (group_ids_in_cluster[1]),
+			    pulongest (group_ids_in_cluster[2]))
+	   : "(?,?,?)");
+
+      grid_pos = string_printf ("%s/%s", grid_pos.c_str (),
+				cluster_pos.c_str ());
+    }
+
+  return grid_pos;
 }
 
 std::string
@@ -516,9 +534,41 @@ wave_coordinates::fetch ()
   amd_dbgapi_wave_get_info (wave_id, AMD_DBGAPI_WAVE_INFO_DISPATCH,
 			    sizeof (dispatch_id), &dispatch_id);
 
+  vec3_u32_t cluster_sizes;
+  if (amd_dbgapi_dispatch_get_info (dispatch_id,
+				    AMD_DBGAPI_DISPATCH_INFO_CLUSTER_SIZES,
+				    sizeof (cluster_sizes), &cluster_sizes)
+      == AMD_DBGAPI_STATUS_SUCCESS)
+    {
+      vec3_t<uint16_t> workgroup_sizes;
+      auto query = AMD_DBGAPI_DISPATCH_INFO_WORKGROUP_SIZES;
+      if (amd_dbgapi_dispatch_get_info (dispatch_id, query,
+					sizeof (workgroup_sizes),
+					&workgroup_sizes)
+	  == AMD_DBGAPI_STATUS_SUCCESS)
+	{
+	  in_cluster = true;
+	  cluster_dims[0] = cluster_sizes[0] / workgroup_sizes[0];
+	  cluster_dims[1] = cluster_sizes[1] / workgroup_sizes[1];
+	  cluster_dims[2] = cluster_sizes[2] / workgroup_sizes[2];
+	}
+    }
+
+  if (!in_cluster)
+    {
+      cluster_dims[0] = 1;
+      cluster_dims[1] = 1;
+      cluster_dims[2] = 1;
+    }
+
+  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_WORKGROUP_COORD,
-			    sizeof (group_ids), &group_ids);
+			    AMD_DBGAPI_WAVE_INFO_WORKGROUP_COORD_IN_CLUSTER,
+			    sizeof (group_ids_in_cluster),
+			    &group_ids_in_cluster);
 
   amd_dbgapi_wave_get_info (wave_id,
 			    AMD_DBGAPI_WAVE_INFO_WAVE_NUMBER_IN_WORKGROUP,
@@ -661,10 +711,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)
@@ -685,25 +735,40 @@ lane_workgroup_pos (thread_info *tp, int lane)
 			   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;
   for (int i = 0; i < 3; i++)
     {
-      size_t work_item_start
-	= static_cast<size_t> (info.coords.group_ids[i]) * work_group_sizes[i];
+      size_t num_wgs = (info.coords.cluster_ids[i]
+			* info.coords.cluster_dims[i]
+			+ info.coords.group_ids_in_cluster[i]);
+      size_t work_item_start = num_wgs * work_group_sizes[i];
       size_t work_item_end = work_item_start + work_group_sizes[i];
       if (work_item_end > grid_sizes[i])
 	work_item_end = grid_sizes[i];
       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.  */
@@ -1127,10 +1192,18 @@ amd_dbgapi_target::workgroup_grid_pos (thread_info *thr)
     return beneath ()->workgroup_grid_pos (thr);
 
   wave_info &info = get_thread_wave_info (thr);
-  if (info.coords.group_ids[0] == UINT32_MAX)
+  if (info.coords.cluster_ids[0] == UINT32_MAX)
     return std::nullopt;
 
-  return info.coords.group_ids;
+  vec3_u32_t pos
+    { info.coords.cluster_ids[0] * info.coords.cluster_dims[0]
+      + info.coords.group_ids_in_cluster[0],
+      info.coords.cluster_ids[1] * info.coords.cluster_dims[1]
+      + info.coords.group_ids_in_cluster[1],
+      info.coords.cluster_ids[2] * info.coords.cluster_dims[2]
+      + info.coords.group_ids_in_cluster[2]};
+
+  return pos;
 }
 
 /* Implementation of target_ops::workgroup_sizes.  */
@@ -1183,6 +1256,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.in_cluster
+      || (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 +1281,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 +4566,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 +4600,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 (
@@ -4563,7 +4661,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 +4672,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");
@@ -4654,6 +4754,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 (

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<std::array<size_t, 3>> partial_workgroup_sizes
+  (thread_info *tp);
+
 #endif /* GDB_AMD_DBGAPI_TARGET_H */

gdb/amdgpu-tdep.c

@@ -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);

gdb/configure

@@ -25337,19 +25337,19 @@ if test "$gdb_require_amd_dbgapi" = true \
   # version of the library.
 
 pkg_failed=no
-{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for amd-dbgapi >= 0.80.0" >&5
-$as_echo_n "checking for amd-dbgapi >= 0.80.0... " >&6; }
+{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for amd-dbgapi >= 0.81.0" >&5
+$as_echo_n "checking for amd-dbgapi >= 0.81.0... " >&6; }
 
 if test -n "$AMD_DBGAPI_CFLAGS"; then
     pkg_cv_AMD_DBGAPI_CFLAGS="$AMD_DBGAPI_CFLAGS"
  elif test -n "$PKG_CONFIG"; then
     if test -n "$PKG_CONFIG" && \
-    { { $as_echo "$as_me:${as_lineno-$LINENO}: \$PKG_CONFIG --exists --print-errors \"amd-dbgapi >= 0.80.0\""; } >&5
-  ($PKG_CONFIG --exists --print-errors "amd-dbgapi >= 0.80.0") 2>&5
+    { { $as_echo "$as_me:${as_lineno-$LINENO}: \$PKG_CONFIG --exists --print-errors \"amd-dbgapi >= 0.81.0\""; } >&5
+  ($PKG_CONFIG --exists --print-errors "amd-dbgapi >= 0.81.0") 2>&5
   ac_status=$?
   $as_echo "$as_me:${as_lineno-$LINENO}: \$? = $ac_status" >&5
   test $ac_status = 0; }; then
-  pkg_cv_AMD_DBGAPI_CFLAGS=`$PKG_CONFIG --cflags "amd-dbgapi >= 0.80.0" 2>/dev/null`
+  pkg_cv_AMD_DBGAPI_CFLAGS=`$PKG_CONFIG --cflags "amd-dbgapi >= 0.81.0" 2>/dev/null`
 		      test "x$?" != "x0" && pkg_failed=yes
 else
   pkg_failed=yes
@@ -25361,12 +25361,12 @@ if test -n "$AMD_DBGAPI_LIBS"; then
     pkg_cv_AMD_DBGAPI_LIBS="$AMD_DBGAPI_LIBS"
  elif test -n "$PKG_CONFIG"; then
     if test -n "$PKG_CONFIG" && \
-    { { $as_echo "$as_me:${as_lineno-$LINENO}: \$PKG_CONFIG --exists --print-errors \"amd-dbgapi >= 0.80.0\""; } >&5
-  ($PKG_CONFIG --exists --print-errors "amd-dbgapi >= 0.80.0") 2>&5
+    { { $as_echo "$as_me:${as_lineno-$LINENO}: \$PKG_CONFIG --exists --print-errors \"amd-dbgapi >= 0.81.0\""; } >&5
+  ($PKG_CONFIG --exists --print-errors "amd-dbgapi >= 0.81.0") 2>&5
   ac_status=$?
   $as_echo "$as_me:${as_lineno-$LINENO}: \$? = $ac_status" >&5
   test $ac_status = 0; }; then
-  pkg_cv_AMD_DBGAPI_LIBS=`$PKG_CONFIG --libs "amd-dbgapi >= 0.80.0" 2>/dev/null`
+  pkg_cv_AMD_DBGAPI_LIBS=`$PKG_CONFIG --libs "amd-dbgapi >= 0.81.0" 2>/dev/null`
 		      test "x$?" != "x0" && pkg_failed=yes
 else
   pkg_failed=yes
@@ -25411,9 +25411,9 @@ else
         _pkg_short_errors_supported=no
 fi
         if test $_pkg_short_errors_supported = yes; then
-	        AMD_DBGAPI_PKG_ERRORS=`$PKG_CONFIG --short-errors --print-errors --cflags --libs "amd-dbgapi >= 0.80.0" 2>&1`
+	        AMD_DBGAPI_PKG_ERRORS=`$PKG_CONFIG --short-errors --print-errors --cflags --libs "amd-dbgapi >= 0.81.0" 2>&1`
         else
-	        AMD_DBGAPI_PKG_ERRORS=`$PKG_CONFIG --print-errors --cflags --libs "amd-dbgapi >= 0.80.0" 2>&1`
+	        AMD_DBGAPI_PKG_ERRORS=`$PKG_CONFIG --print-errors --cflags --libs "amd-dbgapi >= 0.81.0" 2>&1`
         fi
 	# Put the nasty error message in config.log where it belongs
 	echo "$AMD_DBGAPI_PKG_ERRORS" >&5

gdb/configure.ac

@@ -327,7 +327,7 @@ if test "$gdb_require_amd_dbgapi" = true \
   # stability until amd-dbgapi hits 1.0, but for convenience, still check for
   # greater or equal that version.  It can be handy when testing with a newer
   # version of the library.
-  PKG_CHECK_MODULES([AMD_DBGAPI], [amd-dbgapi >= 0.80.0],
+  PKG_CHECK_MODULES([AMD_DBGAPI], [amd-dbgapi >= 0.81.0],
 		    [has_amd_dbgapi=yes], [has_amd_dbgapi=no])
 
   if test "$has_amd_dbgapi" = "yes"; then