masahi commented on a change in pull request #7172:
URL: https://github.com/apache/tvm/pull/7172#discussion_r549894740
##########
File path: python/tvm/topi/cuda/nms.py
##########
@@ -512,26 +512,50 @@ def calculate_overlap(out_tensor, box_a_idx, box_b_idx):
with ib.new_scope():
nthread_by = batch_size
+ nthread_tx = max_threads
+
by = te.thread_axis("blockIdx.y")
+ tx = te.thread_axis("threadIdx.x")
ib.scope_attr(by, "thread_extent", nthread_by)
+ ib.scope_attr(tx, "thread_extent", nthread_tx)
+
i = by
+
base_idx = i * num_anchors * box_data_length
num_valid_boxes_local = ib.allocate(
"int32", (1,), name="num_valid_boxes_local", scope="local"
)
num_valid_boxes_local[0] = 0
+ nkeep = if_then_else(tvm.tir.all(top_k > 0, top_k < valid_count[i]),
top_k, valid_count[i])
def nms_inner_loop(ib, j):
+ # The box j is valid, invalidate other boxes that overlap with j
above iou_threshold
+
+ # When return_indices is False, no need to populate box_indices
+ if return_indices:
+ orig_idx = sorted_index[i * num_anchors + j]
+ box_indices[i, num_valid_boxes_local[0]] = indices[i, orig_idx]
+
+ # TODO(masahi): Want to do this instead of above, but the
following is eliminated
+ # during codegen
+ # # Only one thread needs to this write
+ # with ib.if_scope(tx == 0):
+ # orig_idx = sorted_index[i * num_anchors + j]
+ # box_indices[i, num_valid_boxes_local[0]] = indices[i,
orig_idx]
+
Review comment:
Sorry, the simplification is done by python:
```
In [2]: from tvm import te
In [3]: te.thread_axis("threadIdx.x")
In [5]: tx == 0
Out[5]: False
```
It seems I can do a little trick to get what I want:
```
In [10]: tx + 0 == 0
Out[10]: <tvm.tir.expr.EqualOp at 0x7fe6fe1b9a00>
```
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
us...@infra.apache.org
