code: def non_max_suppression(
      data,
      valid_count,
      indices,
      max_output_size=-1,
      iou_threshold=0.5,
      force_suppress=False,
      top_k=-1,
      coord_start=2,
      score_index=1,
      id_index=0,
      return_indices=True,
      invalid_to_bottom=False,
  ):
      """Non-maximum suppression operator for object detection.
  
      Parameters
      ----------
      data : tvm.te.Tensor
          3-D tensor with shape [batch_size, num_anchors, elem_length].
          The last dimension should be in format of
          [class_id, score, box_left, box_top, box_right, box_bottom].
          It could be the second output out_tensor of get_valid_counts.
  
      valid_count : tvm.te.Tensor
          1-D tensor for valid number of boxes. It could be the output
          valid_count of get_valid_counts.
  
      indices : tvm.te.Tensor
          2-D tensor with shape [batch_size, num_anchors], represents
          the index of box in original data. It could be the third
          output out_indices of get_valid_counts. The values in the
          second dimension are like the output of arange(num_anchors)
          if get_valid_counts is not used before non_max_suppression.
  
      max_output_size : optional, tvm.te.Tensor or int
          Max number of output valid boxes for each instance.
          By default all valid boxes are returned.
  
      iou_threshold : optional, tvm.te.Tensor or float
          Non-maximum suppression threshold.
  
      force_suppress : optional, boolean
          Whether to suppress all detections regardless of class_id.
  
      top_k : optional, int
          Keep maximum top k detections before nms, -1 for no limit.
  
      coord_start : required, int
          Start index of the consecutive 4 coordinates.
  
      score_index : optional, int
          Index of the scores/confidence of boxes.
  
      id_index : optional, int
          index of the class categories, -1 to disable.
  
      return_indices : boolean
          Whether to return box indices in input data.
  
      invalid_to_bottom : optional, boolean
          Whether to move all valid bounding boxes to the top.
  
      Returns
      -------
      out : tvm.te.Tensor
          3-D tensor with shape [batch_size, num_anchors, elem_length].
  
      Example
      --------
      .. code-block:: python
  
          # An example to use nms
          dshape = (1, 5, 6)
          data = te.placeholder(dshape, name="data")
          valid_count = te.placeholder((dshape[0],), dtype="int32", name="valid_count")
          iou_threshold = 0.7
          force_suppress = True
          top_k = -1
          out = non_max_suppression(data=data, valid_count=valid_count, iou_threshold=iou_threshold,
                                   force_suppress=force_supress, top_k=top_k, return_indices=False)
          np_data = np.random.uniform(dshape)
          np_valid_count = np.array([4])
          s = topi.generic.schedule_nms(out)
          f = tvm.build(s, [data, valid_count, out], "cuda")
          dev = tvm.cuda(0)
          tvm_data = tvm.nd.array(np_data, dev)
          tvm_valid_count = tvm.nd.array(np_valid_count, dev)
          tvm_out = tvm.nd.array(np.zeros(dshape, dtype=data.dtype), dev)
          f(tvm_data, tvm_valid_count, tvm_out)
      """
      data_buf = tvm.tir.decl_buffer(
          data.shape, data.dtype, "data_buf", data_alignment=8
      )
  
      sort_tensor = _get_sorted_indices(
          data, data_buf, score_index, (data.shape[0], data.shape[1])
      )
  
      (
          out_bboxes,
          out_scores,
          out_class_ids,
          out_features,
          box_indices,
          num_valid_boxes,
      ) = _run_nms(
          data,
          data_buf,
          sort_tensor,
          valid_count,
          indices,
          max_output_size,
          iou_threshold,
          force_suppress,
          top_k,
          coord_start,
          id_index,
          score_index,
          return_indices,
      )
  
      if return_indices:
          return [box_indices, num_valid_boxes]
  
      return _concatenate_outputs(
          out_bboxes,
          out_scores,
          out_class_ids,
          out_features,
          data.shape,
          coord_start,
          score_index,
          id_index,
      )
  
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 14 for calling a global function with example usage: non_max_suppression(data, valid_count, indices, max_output_size, iou_threshold, force_suppress, top_k, coord_start, score_index, id_index, return_indices, invalid_to_bottom) and returns: _concatenate_outputs, list