masahi commented on a change in pull request #7172:
URL: https://github.com/apache/tvm/pull/7172#discussion_r549893296
##########
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:
if I instead does `if_scope(tx != 0):`, the body of `if_scope` appears
in the generated code without guarded by if. I definitely believe that the
simpilfier is making a wrong decision that `tx` never equals 0 here @tqchen
----------------------------------------------------------------
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:
[email protected]
