You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by ek...@apache.org on 2023/07/31 15:23:33 UTC
[tvm] branch main updated: [FRONTEND][TFLITE] Add support for TFLite's regular NMS operator (#15117)
This is an automated email from the ASF dual-hosted git repository.
ekalda pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm.git
The following commit(s) were added to refs/heads/main by this push:
new 619bb1d939 [FRONTEND][TFLITE] Add support for TFLite's regular NMS operator (#15117)
619bb1d939 is described below
commit 619bb1d939588cd5972ccef64eb120701dceb5c6
Author: Ilya Gozman <92...@users.noreply.github.com>
AuthorDate: Mon Jul 31 19:23:25 2023 +0400
[FRONTEND][TFLITE] Add support for TFLite's regular NMS operator (#15117)
This PR adds support of regular NMS operator.
Open questions:
1. How to properly test added functionality?
Other NMS implementations, e.g., fast NMS, use a TF frozen graph from TF official website to convert a model to TFLite and keep NMS operations only. In order to create a similar test, we need to find an archive on TF official website that contains a frozen graph of a model compiled with --use-regular-nms=True flag. We haven't found it yet, so any help is appreciated.
2. Regular NMS requires two sort operations:
Sorting the scores after selecting scores above nms_score_threshold. This PR implements this with a simple bubble sort in order to prove the algorithm's semantics. We tried to replace it with tvm.contrib.sort.argsort. It works well when testing the regular NMS with run_tvm_graph as fast NMS test does or building and running the regular NMS with llvm target. At the same time, it fails to build (error is provided below) when target=ethos-u,cmsis-nn,c. It seems that __tvm_module_ [...]
The error:
error: '__tvm_module_ctx' undeclared (first use in this function) 203 | if (TVMBackendGetFuncFromEnv(__tvm_module_ctx, "tvm.contrib.sort.argsort", &tvm_contrib_sort_argsort_packed) != 0) {
Sorting the scores of previous and current NMS steps. There are two alternatives here:
implement some sorting algorithm as part of hybrid script (to replace current bubble sort)
save the result of each NMS step and use argsort after the hybrid script part. This approach has a drawback as it requires significant amount of memory to store the results of each NMS step.
---
include/tvm/relay/attrs/vision.h | 25 +++
python/tvm/relay/frontend/tflite.py | 45 ++--
python/tvm/relay/op/strategy/generic.py | 34 ++-
python/tvm/relay/op/vision/_vision.py | 30 +++
python/tvm/relay/op/vision/multibox.py | 22 +-
python/tvm/relay/op/vision/nms.py | 59 +++++
python/tvm/topi/cuda/ssd/multibox.py | 63 +++++-
python/tvm/topi/vision/nms.py | 308 +++++++++++++++++++++++++++
python/tvm/topi/vision/ssd/multibox.py | 47 ++--
src/relay/op/vision/multibox_op.cc | 3 +-
src/relay/op/vision/nms.cc | 61 ++++++
tests/python/frontend/tflite/test_forward.py | 15 ++
tests/python/relay/test_op_level5.py | 31 ++-
tests/scripts/request_hook/request_hook.py | 1 +
14 files changed, 695 insertions(+), 49 deletions(-)
diff --git a/include/tvm/relay/attrs/vision.h b/include/tvm/relay/attrs/vision.h
index 976304e79c..bdb9bfbb49 100644
--- a/include/tvm/relay/attrs/vision.h
+++ b/include/tvm/relay/attrs/vision.h
@@ -61,6 +61,7 @@ struct MultiBoxTransformLocAttrs : public tvm::AttrsNode<MultiBoxTransformLocAtt
bool clip;
double threshold;
Array<IndexExpr> variances;
+ bool keep_background;
TVM_DECLARE_ATTRS(MultiBoxTransformLocAttrs, "relay.attrs.MultiBoxTransformLocAttrs") {
TVM_ATTR_FIELD(clip).set_default(true).describe("Clip out-of-boundary boxes.");
@@ -68,6 +69,9 @@ struct MultiBoxTransformLocAttrs : public tvm::AttrsNode<MultiBoxTransformLocAtt
TVM_ATTR_FIELD(variances)
.set_default(Array<IndexExpr>({0.1f, 0.1f, 0.2f, 0.2f}))
.describe("Variances to be decoded from box regression output.");
+ TVM_ATTR_FIELD(keep_background)
+ .set_default(false)
+ .describe("Whether to keep boxes detected as background or not");
}
};
@@ -129,6 +133,27 @@ struct AllClassNonMaximumSuppressionAttrs
}
};
+/*! \brief Attributes used in regular_non_maximum_suppression operator */
+struct RegularNonMaximumSuppressionAttrs
+ : public tvm::AttrsNode<RegularNonMaximumSuppressionAttrs> {
+ int32_t max_detections_per_class;
+ int32_t max_detections;
+ int32_t num_classes;
+ double iou_threshold;
+ double score_threshold;
+
+ TVM_DECLARE_ATTRS(RegularNonMaximumSuppressionAttrs,
+ "relay.attrs.RegularNonMaximumSuppressionAttrs") {
+ TVM_ATTR_FIELD(max_detections_per_class)
+ .describe("The maxinum number of output selected boxes per class.");
+ TVM_ATTR_FIELD(max_detections).describe("The maxinum number of output selected boxes.");
+ TVM_ATTR_FIELD(num_classes).describe("The number of classes without background.");
+ TVM_ATTR_FIELD(iou_threshold).describe("The IoU threshold for box the overlap test.");
+ TVM_ATTR_FIELD(score_threshold)
+ .describe("Score threshold to filter out low score boxes early.");
+ }
+};
+
/*! \brief Attributes used in roi_align operators */
struct ROIAlignAttrs : public tvm::AttrsNode<ROIAlignAttrs> {
Array<IndexExpr> pooled_size;
diff --git a/python/tvm/relay/frontend/tflite.py b/python/tvm/relay/frontend/tflite.py
index dfc7ed27a4..05ff630353 100644
--- a/python/tvm/relay/frontend/tflite.py
+++ b/python/tvm/relay/frontend/tflite.py
@@ -3443,12 +3443,7 @@ class OperatorConverter(object):
flexbuffer = op.CustomOptionsAsNumpy().tobytes()
custom_options = FlexBufferDecoder(flexbuffer).decode()
- if "use_regular_nms" in custom_options:
- if custom_options["use_regular_nms"]:
- raise tvm.error.OpAttributeUnImplemented(
- "use_regular_nms=True is not yet supported for operator "
- "TFLite_Detection_PostProcess."
- )
+ use_regular_nms = "use_regular_nms" in custom_options and custom_options["use_regular_nms"]
inputs = self.get_input_tensors(op)
assert len(inputs) == 3, "inputs length should be 3"
@@ -3481,15 +3476,14 @@ class OperatorConverter(object):
input_zero_point=inputs[2].qnn_params["zero_point"],
)
- # reshape the cls_pred and loc_prob tensors so
- # they can be consumed by multibox_transform_loc
- cls_pred = _op.transpose(cls_pred, [0, 2, 1])
# loc_prob coords are in yxhw format
# need to convert to xywh
loc_coords = _op.split(loc_prob, 4, axis=2)
loc_prob = _op.concatenate(
[loc_coords[1], loc_coords[0], loc_coords[3], loc_coords[2]], axis=2
)
+ # reshape loc_prob tensor so is can be consumed by
+ # multibox_transform_loc
loc_prob = _op.reshape(loc_prob, [batch_size, anchor_boxes * 4])
# anchor coords are in yxhw format
@@ -3511,13 +3505,41 @@ class OperatorConverter(object):
# attributes for multibox_transform_loc
multibox_transform_loc_attrs = {}
multibox_transform_loc_attrs["clip"] = False
- multibox_transform_loc_attrs["threshold"] = custom_options["nms_score_threshold"]
+ multibox_transform_loc_attrs["threshold"] = (
+ 0.0 if use_regular_nms else custom_options["nms_score_threshold"]
+ )
multibox_transform_loc_attrs["variances"] = (
1 / custom_options["x_scale"],
1 / custom_options["y_scale"],
1 / custom_options["w_scale"],
1 / custom_options["h_scale"],
)
+ multibox_transform_loc_attrs["keep_background"] = use_regular_nms
+
+ ret = _op.vision.multibox_transform_loc(
+ # reshape cls_pred so it can be consumed by
+ # multibox_transform_loc
+ _op.transpose(cls_pred, [0, 2, 1]),
+ loc_prob,
+ anchor_expr,
+ **multibox_transform_loc_attrs,
+ )
+
+ if use_regular_nms:
+ # box coordinates need to be converted from ltrb to (ymin, xmin, ymax, xmax)
+ _, transformed_boxes = _op.split(ret[0], (2,), axis=2)
+ box_l, box_t, box_r, box_b = _op.split(transformed_boxes, 4, axis=2)
+ transformed_boxes = _op.concatenate([box_t, box_l, box_b, box_r], axis=2)
+
+ return _op.vision.regular_non_max_suppression(
+ boxes=transformed_boxes,
+ scores=cls_pred,
+ max_detections_per_class=custom_options["detections_per_class"],
+ max_detections=custom_options["max_detections"],
+ num_classes=custom_options["num_classes"],
+ iou_threshold=custom_options["nms_iou_threshold"],
+ score_threshold=custom_options["nms_score_threshold"],
+ )
# attributes for non_max_suppression
non_max_suppression_attrs = {}
@@ -3528,9 +3550,6 @@ class OperatorConverter(object):
non_max_suppression_attrs["max_output_size"] = custom_options["max_detections"]
non_max_suppression_attrs["invalid_to_bottom"] = False
- ret = _op.vision.multibox_transform_loc(
- cls_pred, loc_prob, anchor_expr, **multibox_transform_loc_attrs
- )
ret = _op.vision.non_max_suppression(ret[0], ret[1], ret[1], **non_max_suppression_attrs)
ret = _op.vision.get_valid_counts(ret, 0)
valid_count = ret[0]
diff --git a/python/tvm/relay/op/strategy/generic.py b/python/tvm/relay/op/strategy/generic.py
index 533d65ead7..ce82c6bd6f 100644
--- a/python/tvm/relay/op/strategy/generic.py
+++ b/python/tvm/relay/op/strategy/generic.py
@@ -1184,7 +1184,10 @@ def wrap_compute_multibox_transform_loc(topi_compute):
clip = bool(get_const_int(attrs.clip))
threshold = get_const_float(attrs.threshold)
variances = get_float_tuple(attrs.variances)
- return topi_compute(inputs[0], inputs[1], inputs[2], clip, threshold, variances)
+ keep_background = bool(get_const_int(attrs.keep_background))
+ return topi_compute(
+ inputs[0], inputs[1], inputs[2], clip, threshold, variances, keep_background
+ )
return _compute_multibox_transform_loc
@@ -1316,6 +1319,35 @@ def all_class_nms_strategy(attrs, inputs, out_type, target):
return strategy
+def wrap_compute_regular_nms(topi_compute):
+ """wrap regular nms topi compute"""
+
+ def _compute_nms(attrs, inputs, out_type):
+ return topi_compute(
+ inputs[0],
+ inputs[1],
+ attrs.max_detections_per_class,
+ attrs.max_detections,
+ attrs.num_classes,
+ attrs.iou_threshold,
+ attrs.score_threshold,
+ )
+
+ return _compute_nms
+
+
+@override_native_generic_func("regular_non_max_suppression_strategy")
+def regular_nms_strategy(attrs, inputs, out_type, target):
+ """regular nms generic strategy"""
+ strategy = _op.OpStrategy()
+ strategy.add_implementation(
+ wrap_compute_regular_nms(topi.vision.regular_non_max_suppression),
+ wrap_topi_schedule(topi.generic.schedule_nms),
+ name="regular_nms.generic",
+ )
+ return strategy
+
+
# roi_align
def wrap_compute_roi_align(topi_compute):
"""wrap roi_align topi compute"""
diff --git a/python/tvm/relay/op/vision/_vision.py b/python/tvm/relay/op/vision/_vision.py
index cab9f703e8..522b79d6b5 100644
--- a/python/tvm/relay/op/vision/_vision.py
+++ b/python/tvm/relay/op/vision/_vision.py
@@ -48,6 +48,9 @@ reg.register_pattern("vision.non_max_suppression", OpPattern.OPAQUE)
reg.register_strategy("vision.all_class_non_max_suppression", strategy.all_class_nms_strategy)
reg.register_pattern("vision.all_class_non_max_suppression", OpPattern.OPAQUE)
+reg.register_strategy("vision.regular_non_max_suppression", strategy.regular_nms_strategy)
+reg.register_pattern("vision.regular_non_max_suppression", OpPattern.OPAQUE)
+
@script
def _get_valid_counts_shape_func(data_shape):
@@ -122,6 +125,33 @@ def all_class_nms_shape_func(attrs, inputs, _):
return _all_class_nms_shape_func_tf(inputs[0], inputs[1])
+@script
+def _regular_nms_shape_func(boxes_shape, scores_shape, attrs):
+ out_boxes_shape = output_tensor((3,), "int64")
+ out_classes_shape = output_tensor((2,), "int64")
+ out_scores_shape = output_tensor((2,), "int64")
+ out_num_detections_shape = output_tensor((1,), "int64")
+
+ out_boxes_shape[0] = boxes_shape[0]
+ out_boxes_shape[1] = int64(attrs.max_detections)
+ out_boxes_shape[2] = int64(4)
+
+ out_classes_shape[0] = boxes_shape[0]
+ out_classes_shape[1] = int64(attrs.max_detections)
+
+ out_scores_shape[0] = boxes_shape[0]
+ out_scores_shape[1] = int64(attrs.max_detections)
+
+ out_num_detections_shape[0] = boxes_shape[0]
+
+ return out_boxes_shape, out_classes_shape, out_scores_shape, out_num_detections_shape
+
+
+@reg.register_shape_func("vision.regular_non_max_suppression", False)
+def regular_nms_shape_func(attrs, inputs, _):
+ return _regular_nms_shape_func(inputs[0], inputs[1], attrs)
+
+
@script
def _roi_align_shape_func_nchw(data_shape, rois_shape, pooled_size):
out = output_tensor((4,), "int64")
diff --git a/python/tvm/relay/op/vision/multibox.py b/python/tvm/relay/op/vision/multibox.py
index c97a8c78cf..898ff4ffed 100644
--- a/python/tvm/relay/op/vision/multibox.py
+++ b/python/tvm/relay/op/vision/multibox.py
@@ -53,7 +53,13 @@ def multibox_prior(
def multibox_transform_loc(
- cls_prob, loc_pred, anchor, clip=True, threshold=0.01, variances=(0.1, 0.1, 0.2, 0.2)
+ cls_prob,
+ loc_pred,
+ anchor,
+ clip=True,
+ threshold=0.01,
+ variances=(0.1, 0.1, 0.2, 0.2),
+ keep_background=False,
):
"""Location transformation for multibox detection
@@ -77,10 +83,22 @@ def multibox_transform_loc(
variances : Tuple of float, optional
variances to be decoded from box regression output.
+ keep_background : boolean, optional
+ Whether to keep boxes detected as background or not.
+
Returns
-------
ret : tuple of tvm.relay.Expr
"""
return expr.TupleWrapper(
- _make.multibox_transform_loc(cls_prob, loc_pred, anchor, clip, threshold, variances), 2
+ _make.multibox_transform_loc(
+ cls_prob,
+ loc_pred,
+ anchor,
+ clip,
+ threshold,
+ variances,
+ keep_background,
+ ),
+ 2,
)
diff --git a/python/tvm/relay/op/vision/nms.py b/python/tvm/relay/op/vision/nms.py
index 8c54075d95..4ce2c44275 100644
--- a/python/tvm/relay/op/vision/nms.py
+++ b/python/tvm/relay/op/vision/nms.py
@@ -226,3 +226,62 @@ def all_class_non_max_suppression(
return expr.TupleWrapper(out, 2)
return expr.TupleWrapper(out, 3)
+
+
+def regular_non_max_suppression(
+ boxes,
+ scores,
+ max_detections_per_class,
+ max_detections,
+ num_classes,
+ iou_threshold,
+ score_threshold,
+):
+ """Regular non-maximum suppression operator for object detection, corresponding to TFLite's
+ regular NMS. NMS is performed for each class separately.
+
+ Parameters
+ ----------
+ boxes : relay.Expr
+ 3-D tensor with shape (batch_size, num_boxes, 4). The four values in boxes
+ encode (ymin, xmin, ymax, xmax) coordinates of a box
+
+ scores: relay.Expr
+ 3-D tensor with shape (batch_size, num_boxes, num_classes_with_background)
+
+ max_detections_per_class : int
+ The maxinum number of output selected boxes per class
+
+ max_detections : int
+ The maxinum number of output selected boxes
+
+ num_classes : int
+ The number of classes without background
+
+ iou_threshold : float
+ IoU test threshold
+
+ score_threshold : float
+ Score threshold to filter out low score boxes early
+
+ Returns
+ -------
+ out : relay.Tuple
+ The output is a relay.Tuple of four tensors. The first is `detection_boxes` of size
+ `(batch_size, max_detections , 4)`, the second is `detection_classes` of size
+ `(batch_size, max_detections)`, the third is `detection_scores` of size
+ `(batch_size, max_detections)`, and the fourth is `num_detections` of size `(batch_size,)`
+ representing the total number of selected boxes per batch.
+ """
+ return expr.TupleWrapper(
+ _make.regular_non_max_suppression(
+ boxes,
+ scores,
+ max_detections_per_class,
+ max_detections,
+ num_classes,
+ iou_threshold,
+ score_threshold,
+ ),
+ 4,
+ )
diff --git a/python/tvm/topi/cuda/ssd/multibox.py b/python/tvm/topi/cuda/ssd/multibox.py
index fba293934d..e401547db4 100644
--- a/python/tvm/topi/cuda/ssd/multibox.py
+++ b/python/tvm/topi/cuda/ssd/multibox.py
@@ -158,7 +158,15 @@ def multibox_prior(data, sizes=(1,), ratios=(1,), steps=(-1, -1), offsets=(0.5,
return out
-def transform_loc_pre(cls_prob, valid_count, temp_valid_count, temp_cls_id, temp_score, threshold):
+def transform_loc_pre(
+ cls_prob,
+ valid_count,
+ temp_valid_count,
+ temp_cls_id,
+ temp_score,
+ threshold,
+ keep_background,
+):
"""Low level IR routing for transform location data preparation.
Parameters
@@ -181,6 +189,9 @@ def transform_loc_pre(cls_prob, valid_count, temp_valid_count, temp_cls_id, temp
threshold : float
Threshold to be a positive prediction.
+ keep_background : int
+ 1 to keep background, 0 to remove it.
+
Returns
-------
stmt : Stmt
@@ -199,6 +210,7 @@ def transform_loc_pre(cls_prob, valid_count, temp_valid_count, temp_cls_id, temp
score = ib.buffer_ptr(temp_score)
threshold = tvm.tir.FloatImm("float32", threshold)
+ keep_background = tvm.tir.IntImm("int8", keep_background)
max_threads = int(tvm.target.Target.current(allow_none=False).max_num_threads)
nthread_tx = max_threads
@@ -217,13 +229,14 @@ def transform_loc_pre(cls_prob, valid_count, temp_valid_count, temp_cls_id, temp
valid_count[i] = 0
score[tid] = -1.0
cls_id[tid] = 0
- with ib.for_range(0, num_classes - 1) as k:
- temp = cls_prob[i * num_classes * num_anchors + (k + 1) * num_anchors + j]
- cls_id[tid] = if_then_else(temp > score[tid], k + 1, cls_id[tid])
- score[tid] = tvm.te.max(temp, score[tid])
+ with ib.for_range(0, num_classes) as k:
+ with ib.if_scope(tvm.tir.any(keep_background == 1, k > 0)):
+ temp = cls_prob[i * num_classes * num_anchors + k * num_anchors + j]
+ cls_id[tid] = if_then_else(temp > score[tid], k, cls_id[tid])
+ score[tid] = tvm.te.max(temp, score[tid])
with ib.if_scope(tvm.tir.all(cls_id[tid] > 0, score[tid] < threshold)):
cls_id[tid] = 0
- with ib.if_scope(cls_id[tid] > 0):
+ with ib.if_scope(tvm.tir.any(keep_background == 1, cls_id[tid] > 0)):
temp_valid_count[tid] = 1
with ib.else_scope():
temp_valid_count[tid] = 0
@@ -250,6 +263,7 @@ def transform_loc_ir(
variances,
batch_size,
num_anchors,
+ keep_background,
):
"""Low level IR routing for transform location in multibox_detection operator.
@@ -285,6 +299,9 @@ def transform_loc_ir(
num_anchors : int
Number of anchors
+ keep_background : int
+ 1 to keep background, 0 to remove it.
+
Returns
-------
stmt : Stmt
@@ -325,6 +342,8 @@ def transform_loc_ir(
score = ib.buffer_ptr(temp_score)
out_loc = ib.buffer_ptr(out)
+ keep_background = tvm.tir.IntImm("int8", keep_background)
+
max_threads = int(tvm.target.Target.current(allow_none=False).max_num_threads)
nthread_tx = max_threads
nthread_bx = (batch_size * num_anchors) // max_threads + 1
@@ -341,10 +360,12 @@ def transform_loc_ir(
i = idxd(tid, num_anchors)
j = idxm(tid, num_anchors)
- with ib.if_scope(cls_id[tid] > 0):
+ with ib.if_scope(tvm.tir.any(keep_background == 1, cls_id[tid] > 0)):
with ib.if_scope(j == 0):
out_base_idx = i * num_anchors * 6
- out_loc[out_base_idx] = cls_id[tid] - 1.0
+ out_loc[out_base_idx] = (
+ cls_id[tid] * 1.0 if keep_background == 1 else cls_id[tid] - 1.0
+ )
out_loc[out_base_idx + 1] = score[tid]
(
out_loc[out_base_idx + 2],
@@ -364,7 +385,9 @@ def transform_loc_ir(
)
with ib.else_scope():
out_base_idx = i * num_anchors * 6 + temp_valid_count[tid - 1] * 6
- out_loc[out_base_idx] = cls_id[tid] - 1.0
+ out_loc[out_base_idx] = (
+ cls_id[tid] * 1.0 if keep_background == 1 else cls_id[tid] - 1.0
+ )
out_loc[out_base_idx + 1] = score[tid]
(
out_loc[out_base_idx + 2],
@@ -387,7 +410,13 @@ def transform_loc_ir(
def multibox_transform_loc(
- cls_prob, loc_pred, anchor, clip=True, threshold=0.01, variances=(0.1, 0.1, 0.2, 0.2)
+ cls_prob,
+ loc_pred,
+ anchor,
+ clip=True,
+ threshold=0.01,
+ variances=(0.1, 0.1, 0.2, 0.2),
+ keep_background=False,
):
"""Location transformation for multibox detection
@@ -411,6 +440,9 @@ def multibox_transform_loc(
variances : tuple of float
Variances to be decoded from box regression output.
+ keep_background : boolean
+ Whether to keep boxes detected as background or not.
+
Returns
-------
ret : tuple of tvm.te.Tensor composed of
@@ -481,7 +513,15 @@ def multibox_transform_loc(
),
],
[cls_prob],
- lambda ins, outs: transform_loc_pre(ins[0], outs[0], outs[1], outs[2], outs[3], threshold),
+ lambda ins, outs: transform_loc_pre(
+ ins[0],
+ outs[0],
+ outs[1],
+ outs[2],
+ outs[3],
+ threshold,
+ int(keep_background),
+ ),
dtype=[valid_count_dtype, valid_count_dtype, valid_count_dtype, cls_prob.dtype],
out_buffers=[valid_count_buf, temp_valid_count_buf, temp_cls_id_buf, temp_score_buf],
tag="multibox_transform_loc_phase_one",
@@ -501,6 +541,7 @@ def multibox_transform_loc(
variances,
batch_size,
num_anchors,
+ int(keep_background),
),
in_buffers=[
loc_pred_buf,
diff --git a/python/tvm/topi/vision/nms.py b/python/tvm/topi/vision/nms.py
index 7a51946a27..7bd94745e2 100644
--- a/python/tvm/topi/vision/nms.py
+++ b/python/tvm/topi/vision/nms.py
@@ -873,3 +873,311 @@ def all_class_non_max_suppression(
)
return [selected_indices, selected_scores, num_total_detections]
+
+
+@hybrid.script
+def hybrid_regular_nms(
+ boxes,
+ scores,
+ max_detections_per_class,
+ max_detections,
+ batch_size,
+ num_boxes,
+ num_classes,
+ num_classes_with_background,
+ iou_threshold,
+ score_threshold,
+):
+ """Hybrid routing for regular non-maximum suppression.
+
+ 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_boxes, num_classes_with_background)
+
+ max_detections_per_class : tvm.tir.const
+ The maxinum number of output selected boxes per class
+
+ max_detections : tvm.tir.const
+ The maxinum number of output selected boxes
+
+ batch_size : tvm.tir.IntImm or tvm.tir.Var
+ The number of batches
+
+ num_boxes : tvm.tir.IntImm or tvm.tir.Var
+ The number of bounding boxes
+
+ num_classes : tvm.tir.const
+ The number of classes without background
+
+ num_classes_with_background : tvm.tir.IntImm or tvm.tir.Var
+ The number of classes including background ones
+
+ iou_threshold : tvm.tir.const
+ IoU test threshold
+
+ score_threshold : tvm.tir.const
+ Score threshold to filter out low score boxes early
+
+ Returns
+ -------
+ detection_boxes : tvm.te.Tensor
+ 3-D tensor with shape [batch_size, max_detections, 4].
+
+ detection_classes : tvm.te.Tensor
+ 2-D tensor with shape [batch_size, max_detections].
+
+ detection_scores : tvm.te.Tensor
+ 2-D tensor with shape [batch_size, max_detections].
+
+ num_detections : tvm.te.Tensor
+ 1-D tensor with shape [batch_size].
+ """
+ # output tensors
+ detection_boxes = output_tensor((batch_size, max_detections, 4), boxes.dtype)
+ detection_classes = output_tensor((batch_size, max_detections), "int32")
+ detection_scores = output_tensor((batch_size, max_detections), scores.dtype)
+ num_detections = output_tensor((batch_size,), "int32")
+
+ # scratch buffers
+ class_scores = allocate((num_boxes,), scores.dtype)
+ keep_indices = allocate((num_boxes,), "int32")
+ keep_scores = allocate((num_boxes,), scores.dtype)
+ sorted_indices = allocate((max_detections + num_boxes,), "int32")
+ sorted_scores = allocate((max_detections + num_boxes,), scores.dtype)
+ active_box_candidate = allocate((num_boxes,), "int32")
+ selected = allocate((num_boxes,), "int32")
+ box_indices_after_regular_nms = allocate((max_detections + num_boxes,), "int32")
+ scores_after_regular_nms = allocate((max_detections + num_boxes,), scores.dtype)
+
+ label_offset = num_classes_with_background - num_classes
+ tmp_idx = 0
+
+ for batch_idx in range(batch_size):
+ size_of_sorted_indices = 0
+
+ for class_id in range(num_classes):
+ for box_id in range(num_boxes):
+ # get scores of boxes corresponding to all anchors for single class
+ class_scores[box_id] = scores[batch_idx, box_id, class_id + label_offset]
+
+ # perform non-maximal suppression on single class
+
+ # select detections above score threshold
+ num_scores_kept = 0
+ for i in range(num_boxes):
+ if class_scores[i] >= score_threshold:
+ keep_scores[num_scores_kept] = class_scores[i]
+ keep_indices[num_scores_kept] = i
+ num_scores_kept += 1
+
+ # iota
+ for i in range(num_scores_kept):
+ sorted_indices[i] = i
+ # decreasing sort of scores
+ for i in range(num_scores_kept):
+ for j in range(num_scores_kept - i - 1):
+ if keep_scores[sorted_indices[j]] < keep_scores[sorted_indices[j + 1]]:
+ tmp_idx = sorted_indices[j]
+ sorted_indices[j] = sorted_indices[j + 1]
+ sorted_indices[j + 1] = tmp_idx
+
+ selected_size = 0
+
+ for i in range(num_scores_kept):
+ active_box_candidate[i] = 1
+
+ num_active_candidate = num_scores_kept
+ for i in range(num_scores_kept):
+ if (
+ num_active_candidate != 0
+ and selected_size < min(num_scores_kept, max_detections_per_class)
+ and active_box_candidate[i] == 1
+ ):
+ selected[selected_size] = keep_indices[sorted_indices[i]]
+ selected_size += 1
+
+ active_box_candidate[i] = 0
+ num_active_candidate -= 1
+
+ for j in range(i + 1, num_scores_kept):
+ if active_box_candidate[j] == 1:
+ # compute IOU
+ i_ymin = boxes[batch_idx, keep_indices[sorted_indices[i]], 0]
+ i_xmin = boxes[batch_idx, keep_indices[sorted_indices[i]], 1]
+ i_ymax = boxes[batch_idx, keep_indices[sorted_indices[i]], 2]
+ i_xmax = boxes[batch_idx, keep_indices[sorted_indices[i]], 3]
+
+ j_ymin = boxes[batch_idx, keep_indices[sorted_indices[j]], 0]
+ j_xmin = boxes[batch_idx, keep_indices[sorted_indices[j]], 1]
+ j_ymax = boxes[batch_idx, keep_indices[sorted_indices[j]], 2]
+ j_xmax = boxes[batch_idx, keep_indices[sorted_indices[j]], 3]
+
+ area_i = (i_ymax - i_ymin) * (i_xmax - i_xmin)
+ area_j = (j_ymax - j_ymin) * (j_xmax - j_xmin)
+
+ iou = 0.0
+ if area_i > 0 and area_j > 0:
+ intersection_ymin = max(i_ymin, j_ymin)
+ intersection_xmin = max(i_xmin, j_xmin)
+ intersection_ymax = min(i_ymax, j_ymax)
+ intersection_xmax = min(i_xmax, j_xmax)
+ intersection_area = max(
+ intersection_ymax - intersection_ymin, 0.0
+ ) * max(intersection_xmax - intersection_xmin, 0.0)
+ iou = intersection_area / (area_i + area_j - intersection_area)
+
+ if iou > iou_threshold:
+ active_box_candidate[j] = 0
+ num_active_candidate -= 1
+
+ # end of non-maximal suppression on single class
+
+ # add selected indices from non-max suppression of boxes in this class
+ output_index = size_of_sorted_indices
+ for i in range(selected_size):
+ selected_index = selected[i]
+
+ box_indices_after_regular_nms[output_index] = (
+ selected_index * num_classes_with_background + class_id + label_offset
+ )
+ scores_after_regular_nms[output_index] = class_scores[selected_index]
+
+ output_index += 1
+
+ # sort the max scores among the selected indices
+ # get the indices for top scores
+ num_indices_to_sort = min(output_index, max_detections)
+
+ # iota
+ for i in range(output_index):
+ sorted_indices[i] = i
+ # deacreasing sort of scores
+ for i in range(output_index):
+ for j in range(output_index - i - 1):
+ if (
+ scores_after_regular_nms[sorted_indices[j]]
+ < scores_after_regular_nms[sorted_indices[j + 1]]
+ ):
+ tmp_idx = sorted_indices[j]
+ sorted_indices[j] = sorted_indices[j + 1]
+ sorted_indices[j + 1] = tmp_idx
+
+ # copy values to temporary vectors
+ for i in range(num_indices_to_sort):
+ sorted_scores[i] = scores_after_regular_nms[sorted_indices[i]]
+ sorted_indices[i] = box_indices_after_regular_nms[sorted_indices[i]]
+
+ # copy scores and indices from temporary vectors
+ for i in range(num_indices_to_sort):
+ box_indices_after_regular_nms[i] = sorted_indices[i]
+ scores_after_regular_nms[i] = sorted_scores[i]
+
+ size_of_sorted_indices = num_indices_to_sort
+
+ # fill output tensors
+ for output_box_index in range(max_detections):
+ box_ymin = 0.0
+ box_xmin = 0.0
+ box_ymax = 0.0
+ box_xmax = 0.0
+ class_idx = 0
+ selected_score = 0.0
+
+ if output_box_index < size_of_sorted_indices:
+ anchor_idx = (
+ box_indices_after_regular_nms[output_box_index] // num_classes_with_background
+ )
+
+ box_ymin = boxes[batch_idx, anchor_idx, 0]
+ box_xmin = boxes[batch_idx, anchor_idx, 1]
+ box_ymax = boxes[batch_idx, anchor_idx, 2]
+ box_xmax = boxes[batch_idx, anchor_idx, 3]
+ class_idx = (
+ box_indices_after_regular_nms[output_box_index]
+ - anchor_idx * num_classes_with_background
+ - label_offset
+ )
+ selected_score = scores_after_regular_nms[output_box_index]
+
+ detection_boxes[batch_idx, output_box_index, 0] = box_ymin
+ detection_boxes[batch_idx, output_box_index, 1] = box_xmin
+ detection_boxes[batch_idx, output_box_index, 2] = box_ymax
+ detection_boxes[batch_idx, output_box_index, 3] = box_xmax
+ detection_classes[batch_idx, output_box_index] = class_idx
+ detection_scores[batch_idx, output_box_index] = selected_score
+
+ num_detections[batch_idx] = size_of_sorted_indices
+
+ return detection_boxes, detection_classes, detection_scores, num_detections
+
+
+def regular_non_max_suppression(
+ boxes,
+ scores,
+ max_detections_per_class,
+ max_detections,
+ num_classes,
+ iou_threshold,
+ score_threshold,
+):
+ """Regular non-maximum suppression operator for object detection, corresponding to TFLite's
+ regular NMS. NMS is performed for each class separately.
+
+ Parameters
+ ----------
+ boxes : tvm.te.Tensor
+ 3-D tensor with shape (batch_size, num_boxes, 4). The four values in boxes
+ encode (ymin, xmin, ymax, xmax) coordinates of a box
+
+ scores: tvm.te.Tensor
+ 3-D tensor with shape (batch_size, num_boxes, num_classes_with_background)
+
+ max_detections_per_class : int
+ The maxinum number of output selected boxes per class
+
+ max_detections : int
+ The maxinum number of output selected boxes
+
+ num_classes : int
+ The number of classes without background
+
+ iou_threshold : float
+ IoU test threshold
+
+ score_threshold : float
+ Score threshold to filter out low score boxes early
+
+ Returns
+ -------
+ out : list of tvm.te.Tensor
+ The output is a list of four tensors. The first is `detection_boxes` of size
+ `(batch_size, max_detections , 4)`, the second is `detection_classes` of size
+ `(batch_size, max_detections)`, the third is `detection_scores` of size
+ `(batch_size, max_detections)`, and the fourth is `num_detections` of size `(batch_size,)`
+ representing the total number of selected boxes per batch.
+ """
+ batch_size, num_boxes, num_classes_with_background = scores.shape
+
+ detection_boxes, detection_classes, detection_scores, num_detections = hybrid_regular_nms(
+ boxes=boxes,
+ scores=scores,
+ max_detections_per_class=tvm.tir.const(max_detections_per_class, dtype="int32"),
+ max_detections=tvm.tir.const(max_detections, dtype="int32"),
+ batch_size=batch_size,
+ num_boxes=num_boxes,
+ num_classes=tvm.tir.const(num_classes, dtype="int32"),
+ num_classes_with_background=num_classes_with_background,
+ iou_threshold=tvm.tir.const(iou_threshold, dtype="float32"),
+ score_threshold=tvm.tir.const(score_threshold, dtype="float32"),
+ )
+
+ return [
+ detection_boxes,
+ cast(detection_classes, dtype="float32"),
+ detection_scores,
+ num_detections,
+ ]
diff --git a/python/tvm/topi/vision/ssd/multibox.py b/python/tvm/topi/vision/ssd/multibox.py
index c1d0d958e2..234bfd7953 100644
--- a/python/tvm/topi/vision/ssd/multibox.py
+++ b/python/tvm/topi/vision/ssd/multibox.py
@@ -137,7 +137,7 @@ def multibox_prior(data, sizes=(1,), ratios=(1,), steps=(-1, -1), offsets=(0.5,
@hybrid.script
-def _hybridy_transform_loc(anchor, pred_loc, variance, clip, batch_idx, anchor_idx):
+def _hybrid_transform_loc(anchor, pred_loc, variance, clip, batch_idx, anchor_idx):
"""Transform prior anchor box to output box through location predictions."""
al = anchor[0, anchor_idx, 0]
at = anchor[0, anchor_idx, 1]
@@ -172,7 +172,15 @@ def _hybridy_transform_loc(anchor, pred_loc, variance, clip, batch_idx, anchor_i
@hybrid.script
-def hybrid_multibox_transform_loc(cls_prob, loc_pred, anchor, clip, threshold, variances):
+def hybrid_multibox_transform_loc(
+ cls_prob,
+ loc_pred,
+ anchor,
+ clip,
+ threshold,
+ variances,
+ keep_background,
+):
"""Hybrid routing for transform location in multibox_detection operator.
Parameters
@@ -195,6 +203,9 @@ def hybrid_multibox_transform_loc(cls_prob, loc_pred, anchor, clip, threshold, v
variances : tvm.nd.NDArray
Variances to be decoded from box regression output.
+ keep_background : tvm.tir.const
+ Whether to keep boxes detected as background or not.
+
Returns
-------
out_loc : tvm.te.Tensor or numpy NDArray
@@ -216,27 +227,28 @@ def hybrid_multibox_transform_loc(cls_prob, loc_pred, anchor, clip, threshold, v
out_loc = output_tensor((batch_size, num_anchors, 6), loc_pred.dtype)
valid_count = output_tensor((batch_size,), "int32")
+ start_cls_idx = 0 if keep_background else 1
+
for i in parallel(batch_size):
valid_count[i] = 0
for j in range(num_anchors):
# Find the predicted class id and probability
score = -1.0
cls_id = 0
- for k in range(num_classes):
- if k > 0:
- temp = cls_prob[i, k, j]
- cls_id = k if temp > score else cls_id
- score = max(temp, score)
+ for k in range(start_cls_idx, num_classes):
+ temp = cls_prob[i, k, j]
+ cls_id = k if temp > score else cls_id
+ score = max(temp, score)
if cls_id > 0 and score < threshold:
cls_id = 0
# [id, prob, xmin, ymin, xmax, ymax]
- # Remove background, restore original id
- if cls_id > 0:
- out_loc[i, valid_count[i], 0] = cls_id - 1.0
+ # Remove background if 'keep_background=False', restore original id
+ if keep_background or cls_id > 0:
+ out_loc[i, valid_count[i], 0] = cls_id - 0.0 if keep_background else cls_id - 1.0
out_loc[i, valid_count[i], 1] = score
for l in range(4):
pred_coord[i, l] = loc_pred[i, j * 4 + l]
- out_coord = _hybridy_transform_loc(anchor, pred_coord, variances, clip, i, j)
+ out_coord = _hybrid_transform_loc(anchor, pred_coord, variances, clip, i, j)
out_loc[i, valid_count[i], 2] = out_coord[0]
out_loc[i, valid_count[i], 3] = out_coord[1]
out_loc[i, valid_count[i], 4] = out_coord[2]
@@ -247,7 +259,13 @@ def hybrid_multibox_transform_loc(cls_prob, loc_pred, anchor, clip, threshold, v
def multibox_transform_loc(
- cls_prob, loc_pred, anchor, clip=True, threshold=0.01, variances=(0.1, 0.1, 0.2, 0.2)
+ cls_prob,
+ loc_pred,
+ anchor,
+ clip=True,
+ threshold=0.01,
+ variances=(0.1, 0.1, 0.2, 0.2),
+ keep_background=False,
):
"""Location transformation for multibox detection
@@ -271,10 +289,14 @@ def multibox_transform_loc(
variances : tuple of float
Variances to be decoded from box regression output.
+ keep_background : boolean
+ Whether to keep boxes detected as background or not.
+
Returns
-------
ret : tuple of tvm.te.Tensor
"""
+
return hybrid_multibox_transform_loc(
cls_prob,
loc_pred,
@@ -282,6 +304,7 @@ def multibox_transform_loc(
tvm.tir.const(clip, "bool"),
tvm.tir.const(threshold, "float32"),
tvm.runtime.convert(variances),
+ tvm.tir.const(keep_background, "bool"),
)
diff --git a/src/relay/op/vision/multibox_op.cc b/src/relay/op/vision/multibox_op.cc
index 17d0a47182..c76316aad4 100644
--- a/src/relay/op/vision/multibox_op.cc
+++ b/src/relay/op/vision/multibox_op.cc
@@ -116,11 +116,12 @@ bool MultiBoxTransformLocRel(const Array<Type>& types, int num_inputs, const Att
}
Expr MakeMultiBoxTransformLoc(Expr cls_prob, Expr loc_pred, Expr anchor, bool clip,
- double threshold, Array<IndexExpr> variances) {
+ double threshold, Array<IndexExpr> variances, bool keep_background) {
auto attrs = make_object<MultiBoxTransformLocAttrs>();
attrs->clip = std::move(clip);
attrs->threshold = std::move(threshold);
attrs->variances = std::move(variances);
+ attrs->keep_background = std::move(keep_background);
static const Op& op = Op::Get("vision.multibox_transform_loc");
return Call(op, {cls_prob, loc_pred, anchor}, Attrs(attrs), {});
}
diff --git a/src/relay/op/vision/nms.cc b/src/relay/op/vision/nms.cc
index 8c33c1648c..24873e468a 100644
--- a/src/relay/op/vision/nms.cc
+++ b/src/relay/op/vision/nms.cc
@@ -210,5 +210,66 @@ RELAY_REGISTER_OP("vision.all_class_non_max_suppression")
.set_support_level(5)
.add_type_rel("AllClassNMS", AllClassNMSRel);
+TVM_REGISTER_NODE_TYPE(RegularNonMaximumSuppressionAttrs);
+
+bool RegularNMSRel(const Array<Type>& types, int num_inputs, const Attrs& attrs,
+ const TypeReporter& reporter) {
+ ICHECK_EQ(types.size(), 3);
+ const auto* boxes = types[0].as<TensorTypeNode>();
+ if (boxes == nullptr) return false;
+ const auto* scores = types[1].as<TensorTypeNode>();
+ if (scores == nullptr) return false;
+
+ const auto& boxes_shape = boxes->shape;
+ const auto& scores_shape = scores->shape;
+ ICHECK_EQ(boxes_shape.size(), 3) << "Input boxes should be 3-D.";
+ ICHECK_EQ(scores_shape.size(), 3) << "Input scores count should be 3-D.";
+
+ IndexExpr num_batches = boxes_shape[0];
+
+ const auto* param = attrs.as<RegularNonMaximumSuppressionAttrs>();
+ CHECK(param);
+
+ std::vector<Type> fields;
+ std::vector<IndexExpr> nmsed_boxes_shape{num_batches, param->max_detections, 4};
+ std::vector<IndexExpr> nmsed_classes_shape{num_batches, param->max_detections};
+ std::vector<IndexExpr> nmsed_scores_shape{num_batches, param->max_detections};
+ std::vector<IndexExpr> nmsed_detections_number_shape{num_batches};
+ fields.push_back(TensorType(nmsed_boxes_shape, DataType::Float(32)));
+ fields.push_back(TensorType(nmsed_classes_shape, DataType::Float(32)));
+ fields.push_back(TensorType(nmsed_scores_shape, DataType::Float(32)));
+ fields.push_back(TensorType(nmsed_detections_number_shape, DataType::Int(32)));
+
+ reporter->Assign(types[2], TupleType(Array<Type>(fields)));
+ return true;
+}
+
+Expr MakeRegularNMS(Expr boxes, Expr scores, int32_t max_detections_per_class,
+ int32_t max_detections, int32_t num_classes, double iou_threshold,
+ double score_threshold) {
+ auto attrs = make_object<RegularNonMaximumSuppressionAttrs>();
+ attrs->max_detections_per_class = max_detections_per_class;
+ attrs->max_detections = max_detections;
+ attrs->num_classes = num_classes;
+ attrs->iou_threshold = iou_threshold;
+ attrs->score_threshold = score_threshold;
+ static const Op& op = Op::Get("vision.regular_non_max_suppression");
+ return Call(op, {boxes, scores}, Attrs(attrs), {});
+}
+
+TVM_REGISTER_GLOBAL("relay.op.vision._make.regular_non_max_suppression")
+ .set_body_typed(MakeRegularNMS);
+
+RELAY_REGISTER_OP("vision.regular_non_max_suppression")
+ .describe(R"doc(TBD)doc" TVM_ADD_FILELINE)
+ .set_num_inputs(2)
+ .add_argument("boxes", "Tensor",
+ "3-D tensor with shape (batch_size, num_boxes, 4). The four values in boxes "
+ "encode (ymin, xmin, ymax, xmax) coordinates of a box.")
+ .add_argument("scores", "Tensor",
+ "3-D tensor with shape (batch_size, num_boxes, num_classes_with_background).")
+ .set_support_level(5)
+ .add_type_rel("RegularNMS", RegularNMSRel);
+
} // namespace relay
} // namespace tvm
diff --git a/tests/python/frontend/tflite/test_forward.py b/tests/python/frontend/tflite/test_forward.py
index 4ea82e5b4c..f00624dc03 100644
--- a/tests/python/frontend/tflite/test_forward.py
+++ b/tests/python/frontend/tflite/test_forward.py
@@ -4566,6 +4566,8 @@ def _test_detection_postprocess(tf_model_file, box_encodings_size, class_predict
def test_detection_postprocess():
"""Detection PostProcess"""
+
+ # Fast-NMS
box_encodings_size = (1, 1917, 4)
class_predictions_size = (1, 1917, 91)
tf_model_file = tf_testing.get_workload_official(
@@ -4575,6 +4577,7 @@ def test_detection_postprocess():
)
_test_detection_postprocess(tf_model_file, box_encodings_size, class_predictions_size)
+ # Fast-NMS
box_encodings_size = (1, 2034, 4)
class_predictions_size = (1, 2034, 91)
tf_model_file = download_testdata(
@@ -4583,6 +4586,18 @@ def test_detection_postprocess():
)
_test_detection_postprocess(tf_model_file, box_encodings_size, class_predictions_size)
+ # Regular NMS
+ box_encodings_size = (1, 1917, 4)
+ class_predictions_size = (1, 1917, 91)
+ tf_model_file = download_testdata(
+ (
+ "https://github.com/Grovety/ModelZoo/raw/52fb82156ae8c8e3f62c7d7caf6867b25261dda4/"
+ "models/object_detection/ssd_mobilenet_v1/tflite_int8/tflite_graph_with_regular_nms.pb"
+ ),
+ "tflite_graph_with_regular_nms.pb",
+ )
+ _test_detection_postprocess(tf_model_file, box_encodings_size, class_predictions_size)
+
#######################################################################
# Custom Converter
diff --git a/tests/python/relay/test_op_level5.py b/tests/python/relay/test_op_level5.py
index db910661ca..4f68256ab5 100644
--- a/tests/python/relay/test_op_level5.py
+++ b/tests/python/relay/test_op_level5.py
@@ -630,7 +630,7 @@ def test_non_max_suppression(executor_kind):
@tvm.testing.uses_gpu
def test_multibox_transform_loc(executor_kind):
- def test_default_value():
+ def test_default_value(keep_background):
num_anchors = 3
num_classes = 3
@@ -644,14 +644,26 @@ def test_multibox_transform_loc(executor_kind):
[[[-0.1, -0.1, 0.1, 0.1], [-0.2, -0.2, 0.2, 0.2], [1.2, 1.2, 1.5, 1.5]]]
).astype("float32")
- expected_np_out = np.array(
- [
+ expected_np_out = (
+ np.array(
[
- [1, 0.69999999, 0, 0, 0.10818365, 0.10008108],
- [0, 0.44999999, 1, 1, 1, 1],
- [0, 0.30000001, 0, 0, 0.22903419, 0.20435292],
+ [
+ [2, 0.69999999, 0, 0, 0.10818365, 0.10008108],
+ [0, 0.49999999, 0, 0, 0.22903419, 0.20435292],
+ [1, 0.44999999, 1, 1, 1, 1],
+ ]
]
- ]
+ )
+ if keep_background
+ else np.array(
+ [
+ [
+ [1, 0.69999999, 0, 0, 0.10818365, 0.10008108],
+ [0, 0.44999999, 1, 1, 1, 1],
+ [0, 0.30000001, 0, 0, 0.22903419, 0.20435292],
+ ]
+ ]
+ )
)
cls_prob = relay.var(
@@ -661,7 +673,7 @@ def test_multibox_transform_loc(executor_kind):
anchors = relay.var("anchors", relay.ty.TensorType((1, num_anchors, 4), "float32"))
mtl = relay.vision.multibox_transform_loc(
- cls_prob=cls_prob, loc_pred=loc_pred, anchor=anchors
+ cls_prob=cls_prob, loc_pred=loc_pred, anchor=anchors, keep_background=keep_background
)
ret = run_infer_type(mtl.astuple())
ref_type = relay.ty.TupleType(
@@ -714,7 +726,8 @@ def test_multibox_transform_loc(executor_kind):
)
assert ret.checked_type == ref_type
- test_default_value()
+ test_default_value(keep_background=False)
+ test_default_value(keep_background=True)
test_threshold()
diff --git a/tests/scripts/request_hook/request_hook.py b/tests/scripts/request_hook/request_hook.py
index dd92a92bc5..c4591116e2 100644
--- a/tests/scripts/request_hook/request_hook.py
+++ b/tests/scripts/request_hook/request_hook.py
@@ -213,6 +213,7 @@ URL_MAP = {
"https://github.com/tlc-pack/web-data/raw/main/testdata/microTVM/data/vww_sample_person.jpg": f"{BASE}/tlc-pack/web-data/testdata/microTVM/data/vww_sample_person.jpg",
"https://github.com/tlc-pack/web-data/raw/main/testdata/microTVM/data/vww_sample_not_person.jpg": f"{BASE}/tlc-pack/web-data/testdata/microTVM/data/vww_sample_not_person.jpg",
"https://github.com/tensorflow/tflite-micro/raw/de8f61a074460e1fa5227d875c95aa303be01240/tensorflow/lite/micro/models/keyword_scrambled.tflite": f"{BASE}/models/tflite/keyword_scrambled_8bit.tflite",
+ "https://github.com/Grovety/ModelZoo/raw/52fb82156ae8c8e3f62c7d7caf6867b25261dda4/models/object_detection/ssd_mobilenet_v1/tflite_int8/tflite_graph_with_regular_nms.pb": f"{BASE}/ssd_mobilenet_v1/tflite_int8/tflite_graph_with_regular_nms.pb",
}