code: def _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,
  ):
      """Run NMS using sorted scores."""
      sort_tensor_buf = tvm.tir.decl_buffer(
          sort_tensor.shape,
          sort_tensor.dtype,
          "sort_tensor_buf",
          data_alignment=8,
      )
  
      valid_count_dtype = "int32"
      valid_count_buf = tvm.tir.decl_buffer(
          valid_count.shape,
          valid_count_dtype,
          "valid_count_buf",
          data_alignment=4,
      )
      indices_buf = tvm.tir.decl_buffer(
          indices.shape, indices.dtype, "indices_buf", data_alignment=8
      )
  
      batch_size = data.shape[0]
      num_anchors = data.shape[1]
      # Number of extra features per box beyond coords, score, and id.
      num_features = (
          data.shape[2] - 6 if id_index >= 0 else data.shape[2] - 5
      )
  
      # output shapes
      bbox_shape = (batch_size, num_anchors, 4)
      score_shape = (batch_size, num_anchors)
      class_id_shape = score_shape
      out_features_shape = (batch_size, num_anchors, num_features)
      box_indices_shape = score_shape
      num_valid_boxes_shape = (batch_size, 1)
  
      return te.extern(
          [
              bbox_shape,
              score_shape,
              class_id_shape,
              out_features_shape,
              box_indices_shape,
              num_valid_boxes_shape,
          ],
          [data, sort_tensor, valid_count, indices],
          lambda ins, outs: nms_ir(
              ins[0],
              ins[1],
              ins[2],
              ins[3],
              outs[0],  # sorted bbox
              outs[1],  # sorted scores
              outs[2],  # sorted class ids
              outs[3],  # sorted box feats
              outs[4],  # box_indices
              outs[5],  # num_valid_boxes
              max_output_size,
              iou_threshold,
              force_suppress,
              top_k,
              coord_start,
              id_index,
              score_index,
              return_indices,
          ),
          dtype=[
              data.dtype,
              "float32",
              "float32",
              "float32",
              "int32",
              "int32",
          ],
          in_buffers=[
              data_buf,
              sort_tensor_buf,
              valid_count_buf,
              indices_buf,
          ],
          name="nms",
          tag="nms",
      )
  
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 12 for calling a global function with example usage: _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) and returns: te