You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by GitBox <gi...@apache.org> on 2020/12/28 10:53:58 UTC
[GitHub] [tvm] masahi opened a new pull request #7172: [TOPI] Parallelize GPU NMS inner loop
masahi opened a new pull request #7172:
URL: https://github.com/apache/tvm/pull/7172
----------------------------------------------------------------
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:
users@infra.apache.org
[GitHub] [tvm] mbrookhart commented on pull request #7172: [TOPI] Parallelize GPU NMS inner loop
Posted by GitBox <gi...@apache.org>.
mbrookhart commented on pull request #7172:
URL: https://github.com/apache/tvm/pull/7172#issuecomment-754276271
Kudos!
----------------------------------------------------------------
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:
users@infra.apache.org
[GitHub] [tvm] Laurawly commented on a change in pull request #7172: [TOPI] Parallelize GPU NMS inner loop
Posted by GitBox <gi...@apache.org>.
Laurawly commented on a change in pull request #7172:
URL: https://github.com/apache/tvm/pull/7172#discussion_r549884392
##########
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:
Not sure why codegen eliminates the if condition. But it looks to me that adding that condition won't make the code faster since it introduces thread divergence.
##########
File path: python/tvm/topi/cuda/nms.py
##########
@@ -575,6 +593,10 @@ def nms_inner_loop(ib, j):
nms_inner_loop(ib, j)
num_valid_boxes[i] = num_valid_boxes_local[0]
+ # TODO(masahi): Want to do this instead of above, but the following is eliminated
+ # during codegen
+ # with ib.if_scope(tx == 0):
+ # num_valid_boxes[i] = num_valid_boxes_local[0]
Review comment:
Same here from my understanding.
----------------------------------------------------------------
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:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a change in pull request #7172: [TOPI] Parallelize GPU NMS inner loop
Posted by GitBox <gi...@apache.org>.
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
----------------------------------------------------------------
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:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a change in pull request #7172: [TOPI] Parallelize GPU NMS inner loop
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #7172:
URL: https://github.com/apache/tvm/pull/7172#discussion_r549325084
##########
File path: python/tvm/topi/cuda/nms.py
##########
@@ -575,6 +593,10 @@ def nms_inner_loop(ib, j):
nms_inner_loop(ib, j)
num_valid_boxes[i] = num_valid_boxes_local[0]
+ # TODO(masahi): Want to do this instead of above, but the following is eliminated
+ # during codegen
+ # with ib.if_scope(tx == 0):
+ # num_valid_boxes[i] = num_valid_boxes_local[0]
Review comment:
Here, I want to generate the code like this:
```
if (threadIdx.x == 0) {
num_valid_boxes[i] = num_valid_boxes_local[0]
}
```
but it seems `if_scope(tx == 0)` is eliminated during codegen. I guess this condition is considered a dead code by the simplifier? How can I generate code like this? @tqchen @Laurawly
----------------------------------------------------------------
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:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a change in pull request #7172: [TOPI] Parallelize GPU NMS inner loop
Posted by GitBox <gi...@apache.org>.
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:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a change in pull request #7172: [TOPI] Parallelize GPU NMS inner loop
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #7172:
URL: https://github.com/apache/tvm/pull/7172#discussion_r549325084
##########
File path: python/tvm/topi/cuda/nms.py
##########
@@ -575,6 +593,10 @@ def nms_inner_loop(ib, j):
nms_inner_loop(ib, j)
num_valid_boxes[i] = num_valid_boxes_local[0]
+ # TODO(masahi): Want to do this instead of above, but the following is eliminated
+ # during codegen
+ # with ib.if_scope(tx == 0):
+ # num_valid_boxes[i] = num_valid_boxes_local[0]
Review comment:
Here, I want to generate the code like this:
```
if (threadIdx.x == 0) {
num_valid_boxes[i] = num_valid_boxes_local[0]
}
```
but it seems this `if_scope` is eliminated during codegen. I guess this condition is considered a dead code by the simplifier? How can I generate code like this? @tqchen @Laurawly
----------------------------------------------------------------
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:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #7172: [TOPI] Parallelize GPU NMS inner loop
Posted by GitBox <gi...@apache.org>.
masahi commented on pull request #7172:
URL: https://github.com/apache/tvm/pull/7172#issuecomment-752381473
thanks @Laurawly
----------------------------------------------------------------
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:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a change in pull request #7172: [TOPI] Parallelize GPU NMS inner loop
Posted by GitBox <gi...@apache.org>.
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:
users@infra.apache.org
[GitHub] [tvm] masahi merged pull request #7172: [TOPI] Parallelize GPU NMS inner loop
Posted by GitBox <gi...@apache.org>.
masahi merged pull request #7172:
URL: https://github.com/apache/tvm/pull/7172
----------------------------------------------------------------
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:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a change in pull request #7172: [TOPI] Parallelize GPU NMS inner loop
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #7172:
URL: https://github.com/apache/tvm/pull/7172#discussion_r549892107
##########
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:
Yes, I was hoping that removing pointless reads and writes of the same memory location by 1024 threads could make it a bit faster. At least it should be no worse.
Not sure if there are many traffic if all threads do the same memory read and write.
----------------------------------------------------------------
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:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a change in pull request #7172: [TOPI] Parallelize GPU NMS inner loop
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #7172:
URL: https://github.com/apache/tvm/pull/7172#discussion_r549902996
##########
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:
This issue was fixed, doing the write by 1 thread makes it 0.1 millisecond faster
----------------------------------------------------------------
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:
users@infra.apache.org