code: def all_class_non_max_suppression(
      boxes,
      scores,
      max_output_boxes_per_class,
      iou_threshold,
      score_threshold,
      output_format="onnx",
  ):
      """Non-maximum suppression operator for object detection, corresponding to ONNX
      NonMaxSuppression and TensorFlow combined_non_max_suppression.
      NMS is performed for each class separately.
  
      Parameters
      ----------
      boxes : tvm.te.Tensor
          3-D tensor with shape (batch_size, num_boxes, 4)
  
      scores: tvm.te.Tensor
          3-D tensor with shape (batch_size, num_classes, num_boxes)
  
      max_output_boxes_per_class : int or tvm.te.Tensor, optional
          The maxinum number of output selected boxes per class
  
      iou_threshold : float or tvm.te.Tensor, optionaIl
          IoU test threshold
  
      score_threshold : float or tvm.te.Tensor, optional
          Score threshold to filter out low score boxes early
  
      output_format : str, optional
          "onnx" or "tensorflow", see below
  
      Returns
      -------
      out : list of tvm.te.Tensor
          If `output_format` is "onnx", the output is two tensors. The first is `indices` of size
          `(batch_size * num_class* num_boxes , 3)` and the second is a scalar tensor
          `num_total_detection` of shape `(1,)` representing the total number of selected
          boxes. The three values in `indices` encode batch, class, and box indices.
          Rows of `indices` are ordered such that selected boxes from batch 0, class 0 come
          first, in descending of scores, followed by boxes from batch 0, class 1 etc. Out of
          `batch_size * num_class* num_boxes` rows of indices, only the first `num_total_detection`
          rows are valid.
  
          If `output_format` is "tensorflow", the output is three tensors, the first
          is `indices` of size `(batch_size, num_class * num_boxes , 2)`, the second is `scores` of
          size `(batch_size, num_class * num_boxes)`, and the third is `num_total_detection` of size
          `(batch_size,)` representing the total number of selected boxes per batch. The two values
          in `indices` encode class and box indices. Of num_class * num_boxes boxes in `indices` at
          batch b, only the first `num_total_detection[b]` entries are valid. The second axis of
          `indices` and `scores` are sorted within each class by box scores, but not across classes.
          So the box indices and scores for the class 0 come first in a sorted order, followed by
          the class 1 etc.
      """
      batch, num_class, num_boxes = scores.shape
  
      scores = reshape(scores, (batch * num_class, num_boxes))
      sorted_scores, sorted_indices = _dispatch_sort(
          scores, ret_type="both"
      )
      valid_count = _get_valid_box_count(sorted_scores, score_threshold)
  
      (
          selected_indices,
          selected_scores,
          num_detections,
      ) = run_all_class_nms(
          boxes,
          sorted_scores,
          sorted_indices,
          valid_count,
          max_output_boxes_per_class,
          iou_threshold,
          _nms_loop,
          return_scores=(output_format == "tensorflow"),
      )
  
      if output_format == "onnx":
          row_offsets, num_total_detections = exclusive_scan(
              num_detections,
              return_reduction=True,
              output_dtype="int64",
          )
          selected_indices = collect_selected_indices(
              num_class,
              selected_indices,
              num_detections,
              row_offsets,
              _collect_selected_indices_ir,
          )
          return [selected_indices, num_total_detections]
  
      num_detections_per_batch = reshape(
          num_detections, (batch, num_class)
      )
      row_offsets, num_total_detections = exclusive_scan(
          num_detections_per_batch,
          return_reduction=True,
          output_dtype="int64",
          axis=1,
      )
  
      (
          selected_indices,
          selected_scores,
      ) = collect_selected_indices_and_scores(
          selected_indices,
          selected_scores,
          num_detections_per_batch,
          row_offsets,
          num_total_detections,
          _collect_selected_indices_and_scores_ir,
      )
  
      return [selected_indices, selected_scores, num_total_detections]
  
imports: import tvm
  from tvm import te
  from tvm.contrib import nvcc
  from tvm.contrib.thrust import can_use_thrust, can_use_rocthrust
  from tvm.ir import register_intrin_lowering
  from tvm.tir import if_then_else
  from .sort import argsort, argsort_thrust
  from .scan import exclusive_scan
  from ..utils import ceil_div
  from ..math import cast
  from ..transform import reshape
  from ..vision.nms_util import calculate_overlap, binary_search, collect_selected_indices, collect_selected_indices_and_scores, run_all_class_nms
label_desc: Use image node_id 18 for calling a global function with example usage: all_class_non_max_suppression(boxes, scores, max_output_boxes_per_class, iou_threshold, score_threshold, output_format) and returns: list, list